gpu.nu

packages/gpu/src/gpu.nu — backend-neutral GPU compute interface.

v0.1.0 ships ONE backend: CUDA (src/cuda.nu). This file is the surface a program codes against — Gpu / GpuKernel / GpuBuffer plus open / compile / alloc / upload / download / launch / sync. The CUDA specifics (driver handles, NVRTC, the void** kernel-param ABI) stay in cuda.nu so a second backend (ROCm/HIP, OpenCL, a CPU fallback) can slot in behind the same names without touching callers.

Handles are small value structs (by-value is fine — they wrap opaque i64 device handles, no owned NURL heap state). Kernel arguments are passed as a Vec i of i64-encoded values built with the gpu_arg_ encoders; gpu_launch lays them out as the void* array CUDA expects.

API

& c @ nurl_f32_to_bits f32 x → i

f32 → its 32-bit IEEE-754 bit pattern (for scalar kernel args). The C param is float, so it must be declared f32 (not f/double) or the double NURL passes lands in the register wrong (alpha → 0).

: ~ i __gpu_backend 0

Selected backend for this process: 0 = CUDA (default), 1 = CPU (host C++). Set once by gpu_open; every op below dispatches on it. A single backend per process is the norm (a program opens one device), so a global is enough and keeps the Gpu / GpuKernel / GpuBuffer value structs unchanged.

@ gpu_backend → i

@ gpu_is_cpu → b

& c @ nurl_static_kernel s name → *u

Backend 2: STATIC — precompiled kernels linked into the binary (a generated kernels_static.c provides the strong nurl_static_kernel; runtime_core.c carries a weak NULL stub so every other link works). No dlopen, no compiler on PATH, no CUDA: the backend for wasm builds and for sealed native binaries. Select with NURL_GPU=static, or from code with ( gpu_force_static ) before gpu_open — the latter is what a wasm entry point uses (calling into libcuda probes from a browser module is not an option).

: ~ i __gpu_force 0

@ gpu_force_static → v

& c @ wgpu_pipeline s name → i

Backend 3: WEBGPU — the kernels run as WGSL compute shaders through the browser's (or Deno's) navigator.gpu, driven by host imports the JS embedder implements (see packages/gpu/web/webgpu.js). The onnx kernels are pre-translated to WGSL and looked up by entry name (like the static backend), so gpu_compile passes only the name. Device memory is a GPUBuffer, addressed by an integer id. gpu_download is asynchronous (mapAsync) — the wasm module must be built with asyncify covering the wgpu_download import. Select from a wasm entry with ( gpu_force_webgpu ) before gpu_open.

& c @ wgpu_alloc i bytes → i

& c @ wgpu_free i id → v

& c @ wgpu_upload i id *u host i bytes → i

& c @ wgpu_download *u host i id i bytes → i

& c @ wgpu_dtod i dst i src i bytes → i

& c @ wgpu_launch i pipeline i total *u args i nargs → i

@ gpu_force_webgpu → v

: Gpu { i ordinal i dev i ctx } // an initialised device + context

── handle types ──────────────────────────────────────────────────

: GpuKernel { i module i func } // a compiled, loaded __global__ fn

: GpuBuffer { i dptr i bytes } // a device-memory allocation

@ gpu_device_count → i

Number of CUDA-capable devices visible to the process.

@ gpu_open i ordinal → Gpu

Initialise the driver, bind device ordinal, and create a context — or, if no CUDA device is available (or NURL_GPU=cpu), select the CPU backend. The returned Gpu is "ok" (ctx != 0) for either backend; a CPU Gpu carries dev = -2, ctx = 1 as its marker.

@ gpu_ok Gpu g → b

@ gpu_name Gpu g → s

Human-readable device name (e.g. "NVIDIA GeForce RTX 4090", or "CPU").

@ gpu_close Gpu g → v

@ gpu_sync Gpu g → i

Block until all submitted work on the context completes. 0 == success. The CPU backend runs kernels synchronously, so there is nothing to await.

@ gpu_compile Gpu g s src s name → GpuKernel

Compile CUDA-C src at runtime (NVRTC) and load entry point name. The kernel must be declared extern "C" __global__. Returns a GpuKernel with func == 0 on a compile/load error (log on stderr).

@ gpu_kernel_ok GpuKernel k → b

@ gpu_kernel_free GpuKernel k → v

@ gpu_host_alloc i bytes → *u

@ gpu_host_free * u buf → v

@ gpu_host_set_f32 * u buf i idx f v → v

@ gpu_host_get_f32 * u buf i idx → f

@ gpu_host_set_i32 * u buf i idx i v → v

@ gpu_host_get_i32 * u buf i idx → i

@ gpu_alloc Gpu g i bytes → GpuBuffer

@ gpu_free GpuBuffer b → v

@ gpu_upload GpuBuffer dst * u host → i

Copy the buffer's worth of bytes host → device. 0 == success.

@ gpu_dtod GpuBuffer dst i src_dptr → i

Device → device copy of dst's worth of bytes from a raw device pointer. 0 == success. Both allocations live in the process's shared (primary) context, so any package's buffer is a valid source.

@ gpu_download * u host GpuBuffer src → i

Copy the buffer's worth of bytes device → host. 0 == success.

@ gpu_arg_buffer GpuBuffer b → i

@ gpu_arg_i32 i v → i

@ gpu_arg_i64 i v → i

@ gpu_arg_f32 f v → i

@ gpu_launch GpuKernel k i grid i block ( Vec i ) args → i

── launch ──────────────────────────────────────────────────────── 1-D launch: grid blocks of block threads. args is the Vec built from gpu_arg_; its contiguous i64 backing IS the argument-value array, so we only build the pointer (void*) layer over it. 0 == success.

@ gpu_grid i n i block → i

Convenience: ceil-divide for picking a grid size from N and block size.

---

cpu.nu

packages/gpu/src/cpu.nu — a CPU backend for packages/gpu.

Runs the very same CUDA-C kernels the CUDA backend does, but on the host: each kernel source is wrapped in a CUDA-compatibility shim (blockIdx / threadIdx / blockDim / gridDim as thread-locals, __global__ etc. as no-op macros) plus a generated grid-loop entry point, compiled to a shared object by the system C++ compiler, dlopen'd, and run — with OpenMP parallelising the grid across cores. gpu_launch's void** argument array is byte-for-byte what CUDA passes, so the generated wrapper just casts each slot per the kernel's parameter type: no change to the neutral surface.

This makes an onnx/YOLOE model runnable with no GPU at all (see gpu.nu's backend fallback). The only host requirement is a C++ compiler on PATH — the CPU analogue of NVRTC.

API

& c @ dlopen s path i flags → *u

& c @ dlsym *u handle s name → *u

& c @ dlclose *u handle → i

& c @ system s cmd → i

& c @ nurl_peek_i32 *u base i idx → i32

& c @ nurl_poke_i32 *u base i idx i32 val → v

& c @ nurl_cpu_launch *u fn *u params i grid i block → v

@ cpu_compile s src s name → *u

Compile CUDA-C src (entry name) to a host shared object and dlopen it. Returns the dlopen handle (0 on failure), analogous to cuda_module_load.

@ cpu_function * u handle → i

The generated grid-loop entry point in a compiled module (0 if absent).

@ cpu_module_free * u handle → v

Do NOT dlclose: a kernel module compiled with -fopenmp keeps libgomp worker threads alive; unloading it out from under them crashes libgomp at thread teardown. Leaving the (small) module mapped for the process lifetime is the standard workaround — it's freed when the process exits anyway.

@ cpu_malloc i bytes → i

── "device" memory — plain host RAM on this backend ────────────── Buffers are f32/i32/i64 arrays (4-byte-aligned, sizes a multiple of 4), so copy word-by-word — avoids declaring libc memcpy (the compiler emits its own memcpy for aggregate copies, and a duplicate FFI declaration collides).

@ cpu_free i ptr → v

@ cpu_htod i dst * u host i bytes → i

@ cpu_dtoh * u host i src i bytes → i

@ cpu_launch i fn i params i grid i block → i

Run a compiled kernel: fn is the __cpu_launch pointer, params the same void** array gpu_launch builds for CUDA. 0 == success.

@ cpu_static_header → s

@ cpu_static_unit s ename s src → String

One kernel: its source (C++-only \extern "C"\ stripped) plus the serial launcher __nurl_sl_<entry>(void** p, long long grid, long long block) mirroring cpu_compile's generated __cpu_launch.

@ cpu_static_registry ( Vec String ) entries → String

The registry: a name→launcher table + the strong nurl_static_kernel.

---

main.nu

packages/gpu/src/main.nu — packages/gpu PoC demo.

Proves GPU compute works from pure NURL: it lists the CUDA devices, compiles two CUDA-C kernels at runtime (NVRTC), runs them on the GPU over the neutral gpu.nu interface, and verifies the results against a CPU reference. No CUDA toolkit calls in the program — just gpu.nu.

gpu # run the demo on device 0 gpu <ordinal> # run on a specific device

Build: nurlpkg install gpu (or ./nurl.sh packages/gpu/src/main.nu)

API

@ k_vadd → s

── kernels (CUDA-C, compiled at runtime) ─────────────────────────

@ k_saxpy → s

@ approx f a f b → b

── helpers ───────────────────────────────────────────────────────

@ say s msg → v

@ demo_vadd Gpu g i n → i

── demo: vector add c = a + b ───────────────────────────────────

@ demo_saxpy Gpu g i n → i

── demo: SAXPY y = alpha*x + y (scalar float + int args) ───────

@ main → i

---

cuda.nu

packages/gpu/src/cuda.nu — the CUDA backend for packages/gpu.

Binds the CUDA Driver API (libcuda, ships with the NVIDIA driver — no toolkit needed) and NVRTC (libnvrtc, runtime CUDA-C → PTX compilation) directly via NURL FFI. NO runtime.c bridge: the & \cuda\` / & \nvrtc\` symbols auto-link through nurl.sh ONLY when a program references them, so a GPU-less target (e.g. MILK-V Duo) never grows a CUDA dependency.

This file is the only place external GPU libraries are imported. The neutral surface in gpu.nu delegates here; a future backend (ROCm/HIP, OpenCL, a CPU fallback) would be a sibling file behind the same gpu.nu.

Handle model: every opaque CUDA handle (CUcontext / CUmodule / CUfunction) and every device address (CUdeviceptr, a u64) is carried as a plain i (i64) in NURL and cast to *u at the FFI boundary. Result codes (CUresult / nvrtcResult) are returned as i; 0 == success.

API

& cuda @ cuInit i32 flags → i32

── CUDA Driver API ───────────────────────────────────────────────

& cuda @ cuDeviceGetCount *u count → i32

& cuda @ cuDeviceGet *u device i32 ordinal → i32

& cuda @ cuDeviceGetName *u name i32 len i32 dev → i32

& cuda @ cuCtxCreate_v2 *u pctx i32 flags i32 dev → i32

& cuda @ cuDevicePrimaryCtxRetain *u pctx i32 dev → i32

& cuda @ cuDevicePrimaryCtxRelease_v2 i32 dev → i32

& cuda @ cuCtxSetCurrent i ctx → i32

& cuda @ cuCtxDestroy_v2 i ctx → i32

& cuda @ cuCtxSynchronize → i32

& cuda @ cuModuleLoadData *u module *u image → i32

& cuda @ cuModuleUnload i module → i32

& cuda @ cuModuleGetFunction *u hfunc i hmod s name → i32

& cuda @ cuMemAlloc_v2 *u dptr i bytesize → i32

& cuda @ cuMemFree_v2 i dptr → i32

& cuda @ cuMemcpyHtoD_v2 i dst *u src i n → i32

& cuda @ cuMemcpyDtoH_v2 *u dst i src i n → i32

& cuda @ cuMemcpyDtoD_v2 i dst i src i n → i32

& cuda @ cuLaunchKernel i f i32 gx i32 gy i32 gz i32 bx i32 by i32 bz i32 sh i stream *u params i extra → i32

& cuda @ cuGetErrorName i32 err *u pstr → i32

& nvrtc @ nvrtcCreateProgram *u prog s src s name i32 nh *u headers *u incs → i32

── NVRTC (runtime CUDA-C → PTX) ────────────────────────────────── nvrtcCreateProgram / nvrtcDestroyProgram take nvrtcProgram (a pointer to the handle slot) → u out-slot offset. The rest take the nvrtcProgram handle BY VALUE → i (i64), portable across native/wasm (see the driver handles above); their remaining *u are genuine out/data offsets.

& nvrtc @ nvrtcCompileProgram i prog i32 nopt *u opts → i32

& nvrtc @ nvrtcGetPTXSize i prog *u sz → i32

& nvrtc @ nvrtcGetPTX i prog *u ptx → i32

& nvrtc @ nvrtcGetProgramLogSize i prog *u sz → i32

& nvrtc @ nvrtcGetProgramLog i prog *u log → i32

& nvrtc @ nvrtcDestroyProgram *u prog → i32

& c @ nurl_peek_f32 *u base i idx → f

── typed 4-byte buffer accessors (runtime.c, always linked) ────── f32 / i32 arrays are the native GPU element layout; the stock 8-byte nurl_peek/poke can't address them at their natural stride.

& c @ nurl_poke_f32 *u base i idx f val → v

& c @ nurl_peek_i32 *u base i idx → i32

nurl_peek_i32 returns a C int and nurl_poke_i32 takes a C int value — both 32-bit. Declaring them with i64 (i) worked on LP64 native (the ABI register ignores the high bits) but is a hard signature mismatch on wasm, where wasm-ld then replaces the caller with an unreachable stub.

& c @ nurl_poke_i32 *u base i idx i32 val → v

@ cuda_init → i

@ cuda_device_count → i

@ cuda_device i ordinal → i

Device ordinal → CUdevice (itself a small int handle). -1 on failure.

@ cuda_device_name i dev → s

CUdevice → device name string in a fresh buffer. The CALLER owns the returned buffer (free with nurl_free after use) — it cannot be a borrow, there is nothing to borrow from.

@ cuda_ctx_create i dev → i

Retain the device's PRIMARY context and make it current → CUcontext handle (0 on failure). Every consumer in the process that opens the same device shares ONE context (and thus one device address space) — the same convention the CUDA runtime API uses — so device pointers can flow between packages (gpukit buffers into an onnx graph, …). cuCtxCreate would give each opener a private context and pointers could not cross.

@ cuda_ctx_destroy_dev i dev → i

Release one retain on the device's primary context (the ctx handle is not destroyed until every retain is released and the process exits).

@ cuda_ctx_destroy i ctx → i

@ cuda_sync → i

@ cuda_compile s src s name → *u

── NVRTC: compile CUDA-C source → PTX text buffer ──────────────── Returns a NUL-terminated PTX buffer (caller passes to cuda_module_load), or 0 on a compile error after printing the NVRTC log to stderr.

@ cuda_module_load * u ptx → i

── module / function ─────────────────────────────────────────────

@ cuda_module_unload i module → i

@ cuda_function i module s name → i

@ cuda_malloc i bytes → i

── device memory ─────────────────────────────────────────────────

@ cuda_free i dptr → i

@ cuda_htod i dptr * u host i bytes → i

@ cuda_dtoh * u host i dptr i bytes → i

@ cuda_dtod i dst i src i bytes → i

@ cuda_launch i func i grid i block * u params → i

── launch ──────────────────────────────────────────────────────── params is a void** array of pointers to the argument values (built by gpu.nu from a Vec of i64-encoded args). 1-D grid/block for now.

@ cuda_error_name i code → s

CUresult code → driver error-name string (e.g. "CUDA_ERROR_INVALID_VALUE"). The returned pointer is the driver's own static string — only the out-slot is ours to release.