hippo

Search:
Group by:
Source   Edit  

Experimental Nim GPU 'macro macros' for generating GPU kernels and helpers to use them. these macros accept a kernel body and produce a macro for easily using the kernel.Nim Library to enable writing CUDA and HIP kernels in Nim All cuda and hip structures and functions are re-exported and can be used

  • There are 3 sets of function prefixes.
  • hippo* prefixed functions are friendly nim interfaces for either HIP or CUDA
    • This is the recommended way to use this library, as it is the most nim-like
    • These functions check for errors and raise them as exceptions
  • hip* prefixed functions are the raw HIP C++ functions
  • cuda* prefixed functions are the raw CUDA C functions

Types

BlockDim {.importcpp: "const __HIP_Coordinates<__HIP_BlockDim>",
           header: "hip/hip_runtime.h".} = object
  x* {.importc: "x".}: uint32_t ## < x
  y* {.importc: "y".}: uint32_t ## < y
  z* {.importc: "z".}: uint32_t ## < z
Source   Edit   
BlockIdx {.importcpp: "const __HIP_Coordinates<__HIP_BlockIdx>",
           header: "hip/hip_runtime.h".} = object
  x* {.importc: "x".}: uint32_t ## < x
  y* {.importc: "y".}: uint32_t ## < y
  z* {.importc: "z".}: uint32_t ## < z
Source   Edit   
ConstCString {.importc: "const char*".} = object
Source   Edit   
Dim3 {.importcpp: "dim3", header: "hip/hip_runtime.h", bycopy.} = object
  x* {.importc: "x".}: uint32_t ## < x
  y* {.importc: "y".}: uint32_t ## < y
  z* {.importc: "z".}: uint32_t ## < z
Source   Edit   
GpuMemory = object
  p*: pointer
Wrapper around gpu memory for automatic cleanup Source   Edit   
GpuRef = ref GpuMemory
Source   Edit   
GridDim {.importcpp: "const __HIP_Coordinates<__HIP_GridDim>",
          header: "hip/hip_runtime.h".} = object
  x* {.importc: "x".}: uint32_t ## < x
  y* {.importc: "y".}: uint32_t ## < y
  z* {.importc: "z".}: uint32_t ## < z
Source   Edit   
hipDeviceProp_t {.importcpp: "hipDeviceProp_t", header: "hip/hip_runtime.h",
                  incompleteStruct.} = object
  deviceOverlap*: cint
  multiProcessorCount*: cint
  sharedMemPerBlock*: csize_t
  warpSize*: cint
Source   Edit   
hipError_t {.importcpp: "hipError_t", header: "hip/hip_runtime.h".} = cint
Source   Edit   
hipEvent_t {.importcpp: "hipEvent_t", header: "hip/hip_runtime.h".} = pointer
Source   Edit   
hipGraph_t {.importcpp: "hipGraph_t", header: "hip/hip_runtime.h".} = pointer
Source   Edit   
hipGraphExec_t {.importcpp: "hipGraphExec_t", header: "hip/hip_runtime.h".} = pointer
Source   Edit   
hipMemcpyKind {.size: 4, header: "hip/hip_runtime.h", importcpp: "hipMemcpyKind".} = enum
  hipMemcpyHostToHost = 0,  ## < Host-to-Host Copy
  hipMemcpyHostToDevice = 1, ## < Host-to-Device Copy
  hipMemcpyDeviceToHost = 2, ## < Device-to-Host Copy
  hipMemcpyDeviceToDevice = 3, ## < Device-to-Device Copy
  hipMemcpyDefault = 4       ## < Runtime will automatically determine copy-kind based on virtual addresses.
Source   Edit   
hipStream_t {.importcpp: "hipStream_t", header: "hip/hip_runtime.h".} = pointer
Source   Edit   
hipStreamCaptureMode {.size: 4, header: "hip/hip_runtime.h",
                       importcpp: "hipStreamCaptureMode".} = enum
  hipStreamCaptureModeGlobal = 0, hipStreamCaptureModeThreadLocal = 1,
  hipStreamCaptureModeRelaxed = 2
Source   Edit   
size_t = uint64
Source   Edit   
ThreadIdx {.importcpp: "const __HIP_Coordinates<__HIP_ThreadIdx>",
            header: "hip/hip_runtime.h".} = object
  x* {.importc: "x".}: uint32_t ## < x
  y* {.importc: "y".}: uint32_t ## < y
  z* {.importc: "z".}: uint32_t ## < z
Source   Edit   
uint8_t = uint8
Source   Edit   
uint16_t = uint16
Source   Edit   
uint32_t = uint32
Source   Edit   
uint64_t = uint64
Source   Edit   
WmmaFloat8 {.importc: "wmma_float8", header: "hippo_wmma.h", bycopy.} = object
Source   Edit   
WmmaHalf16 {.importc: "wmma_half16", header: "hippo_wmma.h", bycopy.} = object
Source   Edit  

Lets

blockDim {.importc, inject, header: "hip/hip_runtime.h".}: BlockDim
Source   Edit   
blockIdx {.importc, inject, header: "hip/hip_runtime.h".}: BlockIdx
Source   Edit   
gridDim {.importc, inject, header: "hip/hip_runtime.h".}: GridDim
Source   Edit   
threadIdx {.importc, inject, header: "hip/hip_runtime.h".}: ThreadIdx
Source   Edit  

Consts

HippoEventDefault = 0'u32
Source   Edit   
HippoMemcpyDefault = hipMemcpyDefault
Source   Edit   
HippoMemcpyDeviceToDevice = hipMemcpyDeviceToDevice
Source   Edit   
HippoMemcpyDeviceToHost = hipMemcpyDeviceToHost
Source   Edit   
HippoMemcpyHostToDevice = hipMemcpyHostToDevice
Source   Edit   
HippoMemcpyHostToHost = hipMemcpyHostToHost
Source   Edit   
HippoPlatform {.strdefine.} = ""
Source   Edit   
HippoRuntime {.strdefine.} = "HIP"
Source   Edit   
HippoWarpSize = 32
Warp/wavefront size for the current backend. Source   Edit   
WarpSize {.intdefine.} = 32
AMD wavefront size. Defaults to 32 (RDNA 3+). Set -d:WarpSize=64 for GCN/CDNA GPUs (e.g. MI250, MI300) which use wave64. Source   Edit  

Procs

proc `$`(self: ConstCString): string {....raises: [], tags: [], forbids: [].}
Source   Edit   
proc `=destroy`(mem: var GpuMemory) {....raises: [], tags: [], forbids: [].}
Automatically free device memory when the object goes out of scope Source   Edit   
proc amdgcnSdot4(a, b: cint; c: cint): cint {.
    importcpp: "__builtin_amdgcn_sudot4(true, #, true, #, #, false)", nodecl,
    ...raises: [], tags: [], forbids: [].}
Source   Edit   
proc blockReduceSum(v: cfloat; shared: ptr cfloat): cfloat {.used, inline,
    exportc, codegenDecl: "__device__ static inline $# $#$#", ...raises: [],
    tags: [], forbids: [].}
Sum v across every thread of the block. shared must point at a caller-provided {.hippoShared.} array of at least (blockDim.x div HippoWarpSize) floats. The total is valid in thread 0 only; other threads hold partial sums. Source   Edit   
proc cos(x: cdouble): cdouble {.header: "hip/hip_runtime.h",
                                importcpp: "cos(@)", ...raises: [], tags: [],
                                forbids: [].}
Source   Edit   
proc cosf(x: cfloat): cfloat {.header: "hip/hip_runtime.h",
                               importcpp: "cosf(@)", ...raises: [], tags: [],
                               forbids: [].}
Source   Edit   
proc coshf(x: cfloat): cfloat {.header: "hip/hip_runtime.h",
                                importcpp: "coshf(@)", ...raises: [], tags: [],
                                forbids: [].}
Source   Edit   
proc dynamicSharedPtr(): pointer {.header: "hip/hip_runtime.h", importcpp: "({ extern __shared__ __align__(16) char __hippo_dyn_smem[]; (void*)__hippo_dyn_smem; })",
                                   nodecl, ...raises: [], tags: [], forbids: [].}
Base pointer of the kernel dynamic shared memory block (extern __shared__). Source   Edit   
proc exp(x: cdouble): cdouble {.header: "hip/hip_runtime.h",
                                importcpp: "exp(@)", ...raises: [], tags: [],
                                forbids: [].}
Source   Edit   
proc expf(x: cfloat): cfloat {.header: "hip/hip_runtime.h",
                               importcpp: "expf(@)", ...raises: [], tags: [],
                               forbids: [].}
Source   Edit   
proc fabsf(x: cfloat): cfloat {.header: "hip/hip_runtime.h",
                                importcpp: "fabsf(@)", ...raises: [], tags: [],
                                forbids: [].}
Source   Edit   
proc floatToHalf(f: cfloat): uint16 {.header: "hip/hip_fp16.h", importcpp: "__half_raw(__float2half(#)).x",
                                      ...raises: [], tags: [], forbids: [].}
Convert float32 to IEEE 754 half-precision (uint16) using HIP hardware intrinsic. Source   Edit   
proc fmaf(a: cfloat; b: cfloat; c: cfloat): cfloat {.
    header: "hip/hip_runtime.h", importcpp: "fmaf(@)", ...raises: [], tags: [],
    forbids: [].}
Source   Edit   
proc fmaxf(a: cfloat; b: cfloat): cfloat {.header: "hip/hip_runtime.h",
    importcpp: "fmaxf(@)", ...raises: [], tags: [], forbids: [].}
Source   Edit   
proc fminf(a: cfloat; b: cfloat): cfloat {.header: "hip/hip_runtime.h",
    importcpp: "fminf(@)", ...raises: [], tags: [], forbids: [].}
Source   Edit   
proc halfToFloat(h: uint16): cfloat {.header: "hip/hip_fp16.h", importcpp: "[&]{ __half_raw r; r.x = (#); return __half2float(r); }()",
                                      ...raises: [], tags: [], forbids: [].}
Convert IEEE 754 half-precision (uint16) to float32 using HIP hardware intrinsic. Source   Edit   
proc handleError(err: hipError_t) {....raises: [Exception], tags: [], forbids: [].}
Source   Edit   
proc hipAtomicAdd(address: ptr float32; val: float32): float32 {.
    header: "hip/hip_runtime.h", importcpp: "atomicAdd(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipAtomicAdd(address: ptr int32; val: int32): int32 {.
    header: "hip/hip_runtime.h", importcpp: "atomicAdd(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipAtomicAdd(address: ptr uint32; val: uint32): uint32 {.
    header: "hip/hip_runtime.h", importcpp: "atomicAdd(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipAtomicAnd(address: ptr int32; val: int32): int32 {.
    header: "hip/hip_runtime.h", importcpp: "atomicAnd(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipAtomicAnd(address: ptr uint32; val: uint32): uint32 {.
    header: "hip/hip_runtime.h", importcpp: "atomicAnd(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipAtomicCAS(address: ptr int32; compare: int32; val: int32): int32 {.
    header: "hip/hip_runtime.h", importcpp: "atomicCAS(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipAtomicCAS(address: ptr uint32; compare: uint32; val: uint32): uint32 {.
    header: "hip/hip_runtime.h", importcpp: "atomicCAS(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipAtomicExch(address: ptr int32; val: int32): int32 {.
    header: "hip/hip_runtime.h", importcpp: "atomicExch(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipAtomicExch(address: ptr uint32; val: uint32): uint32 {.
    header: "hip/hip_runtime.h", importcpp: "atomicExch(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipAtomicMax(address: ptr int32; val: int32): int32 {.
    header: "hip/hip_runtime.h", importcpp: "atomicMax(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipAtomicMax(address: ptr uint32; val: uint32): uint32 {.
    header: "hip/hip_runtime.h", importcpp: "atomicMax(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipAtomicMin(address: ptr int32; val: int32): int32 {.
    header: "hip/hip_runtime.h", importcpp: "atomicMin(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipAtomicMin(address: ptr uint32; val: uint32): uint32 {.
    header: "hip/hip_runtime.h", importcpp: "atomicMin(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipAtomicOr(address: ptr int32; val: int32): int32 {.
    header: "hip/hip_runtime.h", importcpp: "atomicOr(@)", ...raises: [], tags: [],
    forbids: [].}
Source   Edit   
proc hipAtomicOr(address: ptr uint32; val: uint32): uint32 {.
    header: "hip/hip_runtime.h", importcpp: "atomicOr(@)", ...raises: [], tags: [],
    forbids: [].}
Source   Edit   
proc hipAtomicSub(address: ptr int32; val: int32): int32 {.
    header: "hip/hip_runtime.h", importcpp: "atomicSub(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipAtomicSub(address: ptr uint32; val: uint32): uint32 {.
    header: "hip/hip_runtime.h", importcpp: "atomicSub(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipAtomicXor(address: ptr int32; val: int32): int32 {.
    header: "hip/hip_runtime.h", importcpp: "atomicXor(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipAtomicXor(address: ptr uint32; val: uint32): uint32 {.
    header: "hip/hip_runtime.h", importcpp: "atomicXor(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipDeviceSynchronize(): hipError_t {.header: "hip/hip_runtime.h",
    importcpp: "hipDeviceSynchronize(@)", ...raises: [], tags: [], forbids: [].}
Source   Edit   
proc hipEventCreate(event: ptr hipEvent_t): hipError_t {.
    header: "hip/hip_runtime.h", importcpp: "hipEventCreate(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipEventCreateWithFlags(event: ptr hipEvent_t; flags: uint32_t): hipError_t {.
    header: "hip/hip_runtime.h", importcpp: "hipEventCreateWithFlags(@)",
    ...raises: [], tags: [], forbids: [].}
Source   Edit   
proc hipEventDestroy(event: hipEvent_t): hipError_t {.
    header: "hip/hip_runtime.h", importcpp: "hipEventDestroy(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipEventElapsedTime(ms: ptr cfloat; start: hipEvent_t; stop: hipEvent_t): hipError_t {.
    header: "hip/hip_runtime.h", importcpp: "hipEventElapsedTime(@)",
    ...raises: [], tags: [], forbids: [].}
Source   Edit   
proc hipEventQuery(event: hipEvent_t): hipError_t {.header: "hip/hip_runtime.h",
    importcpp: "hipEventQuery(@)", ...raises: [], tags: [], forbids: [].}
Source   Edit   
proc hipEventRecord(event: hipEvent_t; stream: hipStream_t = nil): hipError_t {.
    header: "hip/hip_runtime.h", importcpp: "hipEventRecord(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipEventSynchronize(event: hipEvent_t): hipError_t {.
    header: "hip/hip_runtime.h", importcpp: "hipEventSynchronize(@)",
    ...raises: [], tags: [], forbids: [].}
Source   Edit   
proc hipFree(ptr: pointer): hipError_t {.header: "hip/hip_runtime.h",
    importcpp: "hipFree(@)", ...raises: [], tags: [], forbids: [].}
Source   Edit   
proc hipGetDevice(device: ptr cint): hipError_t {.header: "hip/hip_runtime.h",
    importcpp: "hipGetDevice(@)", ...raises: [], tags: [], forbids: [].}
Source   Edit   
proc hipGetDeviceProperties(prop: ptr hipDeviceProp_t; device: cint): hipError_t {.
    header: "hip/hip_runtime.h", importcpp: "hipGetDeviceProperties(@)",
    ...raises: [], tags: [], forbids: [].}
Source   Edit   
proc hipGetErrorString(err: hipError_t): ConstCString {.
    header: "hip/hip_runtime.h", importcpp: "hipGetErrorString(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipGetLastError(): hipError_t {.header: "hip/hip_runtime.h",
                                     importcpp: "hipGetLastError()", ...raises: [],
                                     tags: [], forbids: [].}
Source   Edit   
proc hipGraphDestroy(graph: hipGraph_t): hipError_t {.
    header: "hip/hip_runtime.h", importcpp: "hipGraphDestroy(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipGraphExecDestroy(graphExec: hipGraphExec_t): hipError_t {.
    header: "hip/hip_runtime.h", importcpp: "hipGraphExecDestroy(@)",
    ...raises: [], tags: [], forbids: [].}
Source   Edit   
proc hipGraphInstantiate(pGraphExec: ptr hipGraphExec_t; graph: hipGraph_t;
                         errNode: pointer = nil; errLog: pointer = nil;
                         logSize: csize_t = 0): hipError_t {.
    header: "hip/hip_runtime.h", importcpp: "hipGraphInstantiate(@)",
    ...raises: [], tags: [], forbids: [].}
Source   Edit   
proc hipGraphLaunch(graphExec: hipGraphExec_t; stream: hipStream_t): hipError_t {.
    header: "hip/hip_runtime.h", importcpp: "hipGraphLaunch(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipHostAlloc(p: ptr pointer; size: csize_t; flags: uint32_t): hipError_t {.
    header: "hip/hip_runtime.h", importcpp: "hipHostAlloc(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipHostFree(p: pointer): hipError_t {.header: "hip/hip_runtime.h",
    importcpp: "hipHostFree(@)", ...raises: [], tags: [], forbids: [].}
Source   Edit   
proc hipLaunchCooperativeKernel(function_address: pointer; gridDim: Dim3;
                                blockDim: Dim3; kernelParams: ptr pointer;
                                sharedMemBytes: csize_t; stream: hipStream_t): hipError_t {.
    importcpp: "hipLaunchCooperativeKernel(@)", header: "hip/hip_runtime.h",
    ...raises: [], tags: [], forbids: [].}
Source   Edit   
proc hipLaunchKernel(function_address: pointer; numBlocks: Dim3;
                     dimBlocks: Dim3; args: ptr pointer): hipError_t {.
    importcpp: "hipLaunchKernel(@)", header: "hip/hip_runtime.h", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipLaunchKernel(function_address: pointer; numBlocks: Dim3;
                     dimBlocks: Dim3; args: ptr pointer;
                     sharedMemBytes: csize_t; stream: hipStream_t): cint {.
    importcpp: "hipLaunchKernel(@)", header: "hip/hip_runtime.h", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipLaunchKernelGGL(function_address: proc; numBlocks: Dim3;
                        dimBlocks: Dim3; sharedMemBytes: uint32_t;
                        stream: hipStream_t) {.
    importcpp: "hipLaunchKernelGGL(@)", header: "hip/hip_runtime.h", varargs,
    ...raises: [], tags: [], forbids: [].}
Source   Edit   
proc hipMalloc(ptr: ptr pointer; size: csize_t): hipError_t {.
    header: "hip/hip_runtime.h", importcpp: "hipMalloc(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipMemcpy(dst: pointer; src: pointer; size: csize_t; kind: hipMemcpyKind): hipError_t {.
    header: "hip/hip_runtime.h", importcpp: "hipMemcpy(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipMemcpyAsync(dst: pointer; src: pointer; sizeBytes: csize_t;
                    kind: hipMemcpyKind; stream: hipStream_t): hipError_t {.
    header: "hip/hip_runtime.h", importcpp: "hipMemcpyAsync(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipMemcpyToSymbol(symbol: pointer; src: pointer; sizeBytes: csize_t;
                       offset: csize_t = 0;
                       kind: hipMemcpyKind = hipMemcpyHostToDevice): hipError_t {.
    header: "hip/hip_runtime.h", importcpp: "hipMemcpyToSymbol(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipMemset(dst: pointer; value: cint; sizeBytes: csize_t): hipError_t {.
    header: "hip/hip_runtime.h", importcpp: "hipMemset(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipMemsetAsync(dst: pointer; value: cint; sizeBytes: csize_t;
                    stream: hipStream_t): hipError_t {.
    header: "hip/hip_runtime.h", importcpp: "hipMemsetAsync(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipOccupancyMaxActiveBlocksPerMultiprocessor(numBlocks: ptr cint;
    f: pointer; blockSize: cint; dynSharedMemPerBlk: csize_t): hipError_t {.
    importcpp: "hipOccupancyMaxActiveBlocksPerMultiprocessor(@)",
    header: "hip/hip_runtime.h", ...raises: [], tags: [], forbids: [].}
Source   Edit   
proc hippoDeviceInfo(): tuple[cuCount, ldsBytes, warpSize: int] {.
    ...raises: [Exception], tags: [], forbids: [].}
Query the current device for its compute unit count, shared memory bytes per block and warp size. The SIMPLE backend reports (1, 0, 1). Source   Edit   
proc hippoMaxActiveBlocksPerSM(kernel: pointer; blockSize: int;
                               sharedMem: int = 0): int {....raises: [Exception],
    tags: [], forbids: [].}
Query how many blocks of the given kernel can run concurrently per SM/CU. Source   Edit   
proc hippoRandUint32(pos: uint32; seed: uint32): uint32 {.inline, exportc,
    codegenDecl: "__device__ __host__ static inline $# $#$#", ...raises: [],
    tags: [], forbids: [].}
Hash pos and seed into a pseudorandom uint32 with SquirrelNoise5. Callable from both host and device code, and stable across backends. Source   Edit   
proc hippoSyncthreads() {.importcpp: "__syncthreads()",
                          header: "hip/hip_runtime.h", ...raises: [], tags: [],
                          forbids: [].}
Source   Edit   
proc hipStreamBeginCapture(stream: hipStream_t; mode: hipStreamCaptureMode): hipError_t {.
    header: "hip/hip_runtime.h", importcpp: "hipStreamBeginCapture(@)",
    ...raises: [], tags: [], forbids: [].}
Source   Edit   
proc hipStreamCreate(stream: ptr hipStream_t): hipError_t {.
    header: "hip/hip_runtime.h", importcpp: "hipStreamCreate(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipStreamDestroy(stream: hipStream_t): hipError_t {.
    header: "hip/hip_runtime.h", importcpp: "hipStreamDestroy(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipStreamEndCapture(stream: hipStream_t; pGraph: ptr hipGraph_t): hipError_t {.
    header: "hip/hip_runtime.h", importcpp: "hipStreamEndCapture(@)",
    ...raises: [], tags: [], forbids: [].}
Source   Edit   
proc hipStreamSynchronize(stream: hipStream_t): hipError_t {.
    header: "hip/hip_runtime.h", importcpp: "hipStreamSynchronize(@)",
    ...raises: [], tags: [], forbids: [].}
Source   Edit   
proc hipStreamWaitEvent(stream: hipStream_t; event: hipEvent_t;
                        flags: uint32_t = 0'u32): hipError_t {.
    header: "hip/hip_runtime.h", importcpp: "hipStreamWaitEvent(@)", ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc hipSymbol[T](sym: var T): pointer {.header: "hip/hip_runtime.h",
    importcpp: "HIP_SYMBOL(@)", ...raises: [], tags: [], forbids: [].}
Source   Edit   
proc hipSyncthreads() {.importcpp: "__syncthreads()",
                        header: "hip/hip_runtime.h", ...raises: [], tags: [],
                        forbids: [].}
Source   Edit   
proc loadU32(p: ptr uint8): uint32 {.header: "hip/hip_runtime.h", importcpp: "[&]{ unsigned int v; __builtin_memcpy(&v, #, 4); return v; }()",
                                     ...raises: [], tags: [], forbids: [].}
Safe uint32 load from a byte pointer (avoids strict aliasing violations). Source   Edit   
proc log(x: cdouble): cdouble {.header: "hip/hip_runtime.h",
                                importcpp: "log(@)", ...raises: [], tags: [],
                                forbids: [].}
Source   Edit   
proc logf(x: cfloat): cfloat {.header: "hip/hip_runtime.h",
                               importcpp: "logf(@)", ...raises: [], tags: [],
                               forbids: [].}
Source   Edit   
proc newDim3(x: uint32_t = 1; y: uint32_t = 1; z: uint32_t = 1): Dim3 {.
    ...raises: [], tags: [], forbids: [].}
Source   Edit   
proc pow(base: cdouble; exp: cdouble): cdouble {.header: "hip/hip_runtime.h",
    importcpp: "pow(@)", ...raises: [], tags: [], forbids: [].}
Source   Edit   
proc powf(base: cfloat; exp: cfloat): cfloat {.header: "hip/hip_runtime.h",
    importcpp: "powf(@)", ...raises: [], tags: [], forbids: [].}
Source   Edit   
proc roundf(x: cfloat): cfloat {.header: "hip/hip_runtime.h",
                                 importcpp: "roundf(@)", ...raises: [], tags: [],
                                 forbids: [].}
Source   Edit   
proc rsqrtf(x: cfloat): cfloat {.header: "hip/hip_runtime.h",
                                 importcpp: "rsqrtf(@)", ...raises: [], tags: [],
                                 forbids: [].}
Fast reciprocal square root for single-precision float. Source   Edit   
proc shfl(val: cfloat; srcLane: cint): cfloat {.header: "hip/hip_runtime.h",
    importcpp: "__shfl(@)", ...raises: [], tags: [], forbids: [].}
Warp shuffle: read float32 from srcLane (broadcast). Source   Edit   
proc shfl(val: cint; srcLane: cint): cint {.header: "hip/hip_runtime.h",
    importcpp: "__shfl(@)", ...raises: [], tags: [], forbids: [].}
Warp shuffle: read int32 from srcLane (broadcast). Source   Edit   
proc shflDown(val: cfloat; delta: cint): cfloat {.header: "hip/hip_runtime.h",
    importcpp: "__shfl_down(@)", ...raises: [], tags: [], forbids: [].}
Warp shuffle down for float32. Source   Edit   
proc shflDown(val: cfloat; delta: cint; width: cint): cfloat {.
    header: "hip/hip_runtime.h", importcpp: "__shfl_down(@)", ...raises: [],
    tags: [], forbids: [].}
Warp shuffle down for float32 within a sub-warp of width lanes. Source   Edit   
proc shflDown(val: cint; delta: cint): cint {.header: "hip/hip_runtime.h",
    importcpp: "__shfl_down(@)", ...raises: [], tags: [], forbids: [].}
Warp shuffle down for int32. Source   Edit   
proc shflDown(val: cint; delta: cint; width: cint): cint {.
    header: "hip/hip_runtime.h", importcpp: "__shfl_down(@)", ...raises: [],
    tags: [], forbids: [].}
Warp shuffle down for int32 within a sub-warp of width lanes. Source   Edit   
proc shflDown(val: cuint; delta: cint): cuint {.header: "hip/hip_runtime.h",
    importcpp: "__shfl_down(@)", ...raises: [], tags: [], forbids: [].}
Warp shuffle down for uint32. Source   Edit   
proc shflDown(val: cuint; delta: cint; width: cint): cuint {.
    header: "hip/hip_runtime.h", importcpp: "__shfl_down(@)", ...raises: [],
    tags: [], forbids: [].}
Warp shuffle down for uint32 within a sub-warp of width lanes. Source   Edit   
proc shflXor(val: cfloat; laneMask: cint): cfloat {.header: "hip/hip_runtime.h",
    importcpp: "__shfl_xor(@)", ...raises: [], tags: [], forbids: [].}
Warp shuffle xor for float32: read the lane whose id is lane xor laneMask. Source   Edit   
proc shflXor(val: cint; laneMask: cint): cint {.header: "hip/hip_runtime.h",
    importcpp: "__shfl_xor(@)", ...raises: [], tags: [], forbids: [].}
Warp shuffle xor for int32: read the lane whose id is lane xor laneMask. Source   Edit   
proc sin(x: cdouble): cdouble {.header: "hip/hip_runtime.h",
                                importcpp: "sin(@)", ...raises: [], tags: [],
                                forbids: [].}
Source   Edit   
proc sinf(x: cfloat): cfloat {.header: "hip/hip_runtime.h",
                               importcpp: "sinf(@)", ...raises: [], tags: [],
                               forbids: [].}
Source   Edit   
proc sqrt(x: cdouble): cdouble {.header: "hip/hip_runtime.h",
                                 importcpp: "sqrt(@)", ...raises: [], tags: [],
                                 forbids: [].}
Source   Edit   
proc sqrtf(x: cfloat): cfloat {.header: "hip/hip_runtime.h",
                                importcpp: "sqrtf(@)", ...raises: [], tags: [],
                                forbids: [].}
Source   Edit   
proc tanhf(x: cfloat): cfloat {.header: "hip/hip_runtime.h",
                                importcpp: "tanhf(@)", ...raises: [], tags: [],
                                forbids: [].}
Source   Edit   
proc vsubss4(a, b: cint): cint {.importcpp: """({
    int result;
    auto ap = (const signed char*)&(#);
    auto bp = (const signed char*)&(#);
    auto rp = (signed char*)&result;
    for (int i = 0; i < 4; i++) {
      int v = (int)ap[i] - (int)bp[i];
      rp[i] = v < -128 ? -128 : (v > 127 ? 127 : v);
    }
    result;
  })""",
                                 nodecl, ...raises: [], tags: [], forbids: [].}
Source   Edit   
proc warpReduceMax(v: cfloat): cfloat {.used, inline, exportc, codegenDecl: "__device__ static inline $# $#$#",
                                        ...raises: [], tags: [], forbids: [].}
Reduce v to its maximum across every lane of the warp. The maximum is valid in lane 0; other lanes hold partial maxima. Source   Edit   
proc warpReduceSum(v: cfloat): cfloat {.used, inline, exportc, codegenDecl: "__device__ static inline $# $#$#",
                                        ...raises: [], tags: [], forbids: [].}
Sum v across every lane of the warp with a shuffle-down ladder. The total is valid in lane 0; other lanes hold partial sums. Source   Edit   
proc wmmaF32_16x16x16_f16_w32(a, b: WmmaHalf16; c: WmmaFloat8): WmmaFloat8 {.
    importc: "__builtin_amdgcn_wmma_f32_16x16x16_f16_w32", nodecl, ...raises: [],
    tags: [], forbids: [].}
Source   Edit   
proc wmmaGetF32(v: WmmaFloat8; idx: cint): cfloat {.importcpp: "#[#]", nodecl,
    ...raises: [], tags: [], forbids: [].}
Source   Edit   
proc wmmaSetF16(v: var WmmaHalf16; idx: cint; val: cushort) {.
    importcpp: "#[#] = ({unsigned short _t = #; *(_Float16*)&_t;})", nodecl,
    ...raises: [], tags: [], forbids: [].}
Source   Edit   
proc wmmaSetF32(v: var WmmaFloat8; idx: cint; val: cfloat) {.
    importcpp: "#[#] = #", nodecl, ...raises: [], tags: [], forbids: [].}
Source   Edit   
proc wmmaZeroF32(): WmmaFloat8 {.importcpp: "(wmma_float8){0,0,0,0,0,0,0,0}",
                                 nodecl, ...raises: [], tags: [], forbids: [].}
Source   Edit  

Converters

converter toConstCString(self: cstring): ConstCString {.
    importc: "(const char*)", noconv, nodecl, ...raises: [], tags: [], forbids: [].}
Source   Edit   
converter toCString(self: ConstCString): cstring {.importc: "(char*)", noconv,
    nodecl, ...raises: [], tags: [], forbids: [].}
Source   Edit  

Macros

macro generateForLoopMacro(name: static[string]; theType: typedesc;
                           body: untyped): untyped
Generate a GPU kernel and a host launcher proc using a custom name. Source   Edit   
macro hippoArgs(args: varargs[untyped]): untyped
Automatically convert varargs for use with CUDA/HIP. CUDA/HIP expects an array of arguments or pointers depending on platform. Source   Edit   
macro hippoConstant(v: untyped): untyped
Declared a variable as __constant__. Constants are read-only globals that are cached on-chip. constants are useful for data that is being read by all threads in a warp at the same time. if each thread in a warp accesses different addresses in constant memory, the accesses are serialized and this may cause a 16x slowdown. eg: const N {.hippoConstant.} = 1024 Source   Edit   
macro hippoDevice(fn: untyped): untyped
Declare fuctions for use on the __device__ (the gpu), to be called by either device or global functions. Source   Edit   
macro hippoDeviceInline(fn: untyped): untyped
Declare a __device__ function that is emitted into every translation unit that calls it. Device symbols do not link across translation units without relocatable device code, so library device helpers must be inline. Source   Edit   
macro hippoGlobal(fn: untyped): untyped
Declare a function as __global__. global functions are called from the host and run on the device. Source   Edit   
macro hippoHost(fn: untyped): untyped
Explicitly declare a function as a __host__ function (cpu side). All functions default to host functions, so this is not required. Source   Edit   
macro hippoHostDevice(fn: untyped): untyped
Declare a function as both __host__ and __device__. This is useful for functions that are usable from either the host and the device. eg: proc add(a: int, b: int) {.hippoHostDevice.} = a + b Source   Edit   
macro hippoHostDeviceInline(fn: untyped): untyped
Declare a __device__ __host__ function that is emitted into every translation unit that calls it, so it links from both host and device code. Source   Edit   
macro hippoShared(v: untyped): untyped
Declared a variable as static shared memory __shared__. Shared memory is shared between threads in the same block. It is faster than global memory, but is limited in size. They are located on-chip. eg: var cache {.hippoShared.}: array[256, float] Source   Edit  

Templates

template hippoAtomicAdd(address: ptr [float32]; val: float32): float32
Atomically add a float32 and return the previous value. Source   Edit   
template hippoAtomicAdd(address: ptr [int32]; val: int32): int32
Atomically add and return the previous value. Source   Edit   
template hippoAtomicAdd(address: ptr [uint32]; val: uint32): uint32
Atomically add and return the previous value. Source   Edit   
template hippoAtomicAnd(address: ptr [int32]; val: int32): int32
Atomically AND and return the previous value. Source   Edit   
template hippoAtomicAnd(address: ptr [uint32]; val: uint32): uint32
Atomically AND and return the previous value. Source   Edit   
template hippoAtomicCAS(address: ptr [int32]; compare: int32; val: int32): int32
Atomically compare-and-swap and return the previous value. Source   Edit   
template hippoAtomicCAS(address: ptr [uint32]; compare: uint32; val: uint32): uint32
Atomically compare-and-swap and return the previous value. Source   Edit   
template hippoAtomicExch(address: ptr [int32]; val: int32): int32
Atomically exchange and return the previous value. Source   Edit   
template hippoAtomicExch(address: ptr [uint32]; val: uint32): uint32
Atomically exchange and return the previous value. Source   Edit   
template hippoAtomicMax(address: ptr [int32]; val: int32): int32
Atomically update with max and return the previous value. Source   Edit   
template hippoAtomicMax(address: ptr [uint32]; val: uint32): uint32
Atomically update with max and return the previous value. Source   Edit   
template hippoAtomicMin(address: ptr [int32]; val: int32): int32
Atomically update with min and return the previous value. Source   Edit   
template hippoAtomicMin(address: ptr [uint32]; val: uint32): uint32
Atomically update with min and return the previous value. Source   Edit   
template hippoAtomicOr(address: ptr [int32]; val: int32): int32
Atomically OR and return the previous value. Source   Edit   
template hippoAtomicOr(address: ptr [uint32]; val: uint32): uint32
Atomically OR and return the previous value. Source   Edit   
template hippoAtomicSub(address: ptr [int32]; val: int32): int32
Atomically subtract and return the previous value. Source   Edit   
template hippoAtomicSub(address: ptr [uint32]; val: uint32): uint32
Atomically subtract and return the previous value. Source   Edit   
template hippoAtomicXor(address: ptr [int32]; val: int32): int32
Atomically XOR and return the previous value. Source   Edit   
template hippoAtomicXor(address: ptr [uint32]; val: uint32): uint32
Atomically XOR and return the previous value. Source   Edit   
template hippoCos(x: cfloat): cfloat
Cosine function for single-precision float Source   Edit   
template hippoCoshf(x: cfloat): cfloat
Hyperbolic cosine for single-precision float. Source   Edit   
template hippoDynamicShared(T: typedesc): ptr UncheckedArray[T:type]
Access the kernel dynamic shared memory block as an UncheckedArray[T]. The size is the sharedMemBytes argument given to hippoLaunchKernel. Source   Edit   
template hippoEventCreate(): HippoEvent
Create a timing event. Source   Edit   
template hippoEventCreateWithFlags(flags: uint32 = 0'u32): HippoEvent
Create a timing event with explicit runtime flags. Source   Edit   
template hippoEventDestroy(event: HippoEvent)
Destroy a timing event. Source   Edit   
template hippoEventElapsedTime(start: HippoEvent; stop: HippoEvent): float32
Get elapsed time between two events. Source   Edit   
template hippoEventQuery(event: HippoEvent): bool
Return true when the event has completed and false when it is still pending. On HIP-CPU, this behaves as a blocking synchronize because hipEventQuery is not declared. Source   Edit   
template hippoEventRecord(event: HippoEvent; stream: HippoStream = nil)
Record an event on a stream. Source   Edit   
template hippoEventSynchronize(event: HippoEvent)
Wait for an event to complete. Source   Edit   
template hippoExp(x: cdouble): cdouble
Exponential function (e^x) for double-precision float Source   Edit   
template hippoExp(x: cfloat): cfloat
Exponential function (e^x) for single-precision float Source   Edit   
template hippoExpf(x: cfloat): cfloat
Exponential function (e^x) for single-precision float. cfloat-only alias for hippoExp that avoids the f32/f64 overload ambiguity. Source   Edit   
template hippoFabsf(x: cfloat): cfloat
Source   Edit   
template hippoFloatToHalf(f: cfloat): uint16
Convert float32 to IEEE 754 half-precision (uint16). Uses hardware intrinsic on HIP/CUDA, software fallback on SIMPLE. Source   Edit   
template hippoFmaf(a: cfloat; b: cfloat; c: cfloat): cfloat
Fused multiply-add (a * b + c) for single-precision float. Source   Edit   
template hippoFmaxf(a: cfloat; b: cfloat): cfloat
Source   Edit   
template hippoFminf(a: cfloat; b: cfloat): cfloat
Minimum of two single-precision floats. Source   Edit   
template hippoFree(p: pointer)
Free memory on the GPU Source   Edit   
template hippoGetDevice(): cint
Get the current device index. Source   Edit   
template hippoGetDeviceProperties(prop: var HippoDeviceProp; device: cint)
Get device properties for the given device. Source   Edit   
template hippoGraphDestroy(graph: hipGraph_t)
Source   Edit   
template hippoGraphLaunch(exec: hipGraphExec_t; stream: HippoStream)
Source   Edit   
template hippoGridSync()
Grid-wide barrier. All blocks must reach this point before any proceed. Only valid inside a kernel launched with hippoLaunchCooperative. Source   Edit   
template hippoHalfToFloat(h: uint16): cfloat
Convert IEEE 754 half-precision (uint16) to float32. Uses hardware intrinsic on HIP/CUDA, software fallback on SIMPLE. Source   Edit   
template hippoHostAlloc(size: int): pointer
Allocate page-locked host memory Source   Edit   
template hippoHostFree(p: pointer)
Free page-locked host memory. Source   Edit   
template hippoLaunchCooperative(kernel: untyped;
                                gridDim: Dim3 = newDim3(1, 1, 1);
                                blockDim: Dim3 = newDim3(1, 1, 1);
                                sharedMemBytes: uint32 = 0;
                                stream: HippoStream = nil; args: untyped)
Launch a kernel using cooperative launch API, enabling grid-wide sync. The kernel can call hippoGridSync() for grid-level barriers. Source   Edit   
template hippoLaunchKernel(kernel: untyped; ## The GPU kernel procedure to launch
                           gridDim: Dim3 = newDim3(1, 1, 1); ## default to a grid of 1 block
                           blockDim: Dim3 = newDim3(1, 1, 1); ## default to 1 thread per block
                           sharedMemBytes: uint32 = 0; ## dynamic shared memory amount to allocate
                           stream: HippoStream = nil; ## Which device stream to run under (defaults to null)
                           args: untyped)
Source   Edit   
template hippoLoadU32(p: ptr uint8): uint32
Safe uint32 load from a byte pointer using __builtin_memcpy. Avoids strict aliasing violations that break castptr uint32 on HIP/CUDA. Source   Edit   
template hippoLog(x: cdouble): cdouble
Natural logarithm for double-precision float Source   Edit   
template hippoLog(x: cfloat): cfloat
Natural logarithm for single-precision float Source   Edit   
template hippoMalloc(size: int): GpuRef
Allocate memory on the GPU and return a GpuRef object. GpuMemory is a wrapper around Gpu allocated pointers. It will automatically free the memory when it goes out of scope. Source   Edit   
template hippoMemcpy(dst: GpuRef; src: GpuRef; size: int; kind: HippoMemcpyKind)
device -> device memory copy Copy memory from src to dst. direction of device and host is determined by kind. Source   Edit   
template hippoMemcpy(dst: GpuRef; src: pointer; size: int; kind: HippoMemcpyKind)
device -> host memory copy Copy memory from src to dst. direction of device and host is determined by kind. Source   Edit   
template hippoMemcpy(dst: pointer; src: GpuRef; size: int; kind: HippoMemcpyKind)
host -> device memory copy Copy memory from src to dst. direction of device and host is determined by kind. Source   Edit   
template hippoMemcpy(dst: pointer; src: pointer; size: int;
                     kind: HippoMemcpyKind)
host -> host memory copy hippoMemcpy is broken out as 4 separate templates to make it easier to work with GpuRef objects Copy memory from src to dst. direction of device and host is determined by kind. Source   Edit   
template hippoMemcpyAsync(dst: pointer; src: pointer; size: int;
                          kind: HippoMemcpyKind; stream: HippoStream)
Asynchronous memory copy on a stream Source   Edit   
template hippoMemcpyToSymbol(symbol: untyped; src: pointer; size: int;
                             offset: int = 0;
                             kind: HippoMemcpyKind = HippoMemcpyHostToDevice)
Copy data from host to a device symbol (eg: __constant__ storage). Source   Edit   
template hippoMemset(dst: GpuRef; value: cint; size: int)
Fill size bytes of the device allocation dst with the byte value. Source   Edit   
template hippoMemset(dst: pointer; value: cint; size: int)
Fill size bytes of device memory at dst with the byte value. Source   Edit   
template hippoMemsetAsync(dst: GpuRef; value: cint; size: int;
                          stream: HippoStream)
Asynchronously fill size bytes of the device allocation dst on a stream. Source   Edit   
template hippoMemsetAsync(dst: pointer; value: cint; size: int;
                          stream: HippoStream)
Asynchronously fill size bytes of device memory at dst on a stream. Source   Edit   
template hippoPow(base: cdouble; exp: cdouble): cdouble
Power function (base^exp) for double-precision float Source   Edit   
template hippoPow(base: cfloat; exp: cfloat): cfloat
Power function (base^exp) for single-precision float Source   Edit   
template hippoRoundf(x: cfloat): cfloat
Source   Edit   
template hippoRsqrtf(x: cfloat): cfloat
Reciprocal square root (1/sqrt(x)) for single-precision float. Source   Edit   
template hippoSdot4(a, b: cint; c: cint): cint
4×int8 dot product: treats a and b as 4 packed signed int8, accumulates into c. Source   Edit   
template hippoShfl(val: cfloat; srcLane: int): cfloat
Warp shuffle: read float32 from srcLane (broadcast). Source   Edit   
template hippoShfl(val: cint; srcLane: int): cint
Warp shuffle: read int32 from srcLane (broadcast). Source   Edit   
template hippoShflDown(val: cfloat; delta: int): cfloat
Warp shuffle down for float32. Returns the value from the lane delta positions below the calling lane within the same warp. Source   Edit   
template hippoShflDown(val: cfloat; delta: int; width: int): cfloat
Warp shuffle down for float32 within a sub-warp of width lanes. Source   Edit   
template hippoShflDown(val: cint; delta: int): cint
Warp shuffle down for int32. Source   Edit   
template hippoShflDown(val: cint; delta: int; width: int): cint
Warp shuffle down for int32 within a sub-warp of width lanes. Source   Edit   
template hippoShflDown(val: cuint; delta: int): cuint
Warp shuffle down for uint32. Source   Edit   
template hippoShflDown(val: cuint; delta: int; width: int): cuint
Warp shuffle down for uint32 within a sub-warp of width lanes. Source   Edit   
template hippoShflXor(val: cfloat; laneMask: int): cfloat
Warp shuffle xor for float32: read the lane whose id is lane xor laneMask. Source   Edit   
template hippoShflXor(val: cint; laneMask: int): cint
Warp shuffle xor for int32: read the lane whose id is lane xor laneMask. Source   Edit   
template hippoSin(x: cfloat): cfloat
Sine function for single-precision float Source   Edit   
template hippoSqrt(x: cfloat): cfloat
Square root function for single-precision float Source   Edit   
template hippoStreamBeginCapture(stream: HippoStream; mode: hipStreamCaptureMode = hipStreamCaptureModeGlobal)
Source   Edit   
template hippoStreamCreate(): HippoStream
Create a new stream Source   Edit   
template hippoStreamDestroy(stream: HippoStream)
Destroy a stream Source   Edit   
template hippoStreamSynchronize(stream: HippoStream)
Synchronize a specific stream Source   Edit   
template hippoStreamWaitEvent(stream: HippoStream; event: pointer;
                              flags: uint32 = 0'u32)
Block stream work until an event is recorded and completed. Source   Edit   
template hippoSynchronize()
Synchronize the device Source   Edit   
template hippoTanhf(x: cfloat): cfloat
Hyperbolic tangent for single-precision float. Source   Edit   
template hippoVsubss4(a, b: cint): cint
Packed int8×4 saturated subtract: resulti = clamp(ai - bi, -128, 127) Source   Edit   
template hippoWmmaGetF32(v: HippoWmmaFloat8; idx: int): cfloat
Source   Edit   
template hippoWmmaSetF16(v: var HippoWmmaHalf16; idx: int; val: uint16)
Source   Edit   
template hippoWmmaSetF32(v: var HippoWmmaFloat8; idx: int; val: cfloat)
Source   Edit  

Made with Nim. Generated: 2026-09-03 20:50:53 UTC