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_best_device → i

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. Which device to open when the caller does not care: $NURL_GPU_DEVICE (an ordinal), else the best one — highest compute capability, ties broken by memory. A box with an old card beside a new one would otherwise silently run everything on whichever the driver enumerates first.

@ gpu_open i ordinal → Gpu

@ 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_timer_new Gpu g → i

── Timers (CUDA backend only; every other backend reports 0) ───────── A begin/end pair of events recorded on the launch stream. Timing a launch from the host measures the launch; these measure the GPU. The handle is 0 when the backend has no events, and every call below is a no-op on a 0 handle, so a caller need not branch on the backend.

@ gpu_timer_mark Gpu g i ev → v

@ gpu_timer_ns Gpu g i start i end → i

Nanoseconds between two marks. Blocks until the end event has completed on the device, so it is a measurement point, not free.

@ gpu_timer_free Gpu g i ev → v

@ gpu_graph_begin Gpu g → b

── CUDA Graphs (CUDA backend only; every other backend reports F/0) ── Capture the launches between begin and end, then replay them all with ONE gpu_graph_launch — same kernels, same argument values, same order, bit-identical results. Callers must not sync between begin and end.

@ gpu_graph_end Gpu g → i

The executable-graph handle, or 0 when capture/instantiation failed.

@ gpu_graph_launch Gpu g i exec → i

@ gpu_graph_free i exec → v

@ gpu_compile Gpu g s src s name → GpuKernel

@ 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

: ~ i __gpu_stage_a 0

── pinned staging for large uploads (CUDA) ─────────────────────────

cuMemcpyHtoD from PAGEABLE memory makes the driver stage every chunk through its own small pinned buffer on one thread — ~5 GB/s, so a 17 GB model spends ~3.5 s of pure CPU in the driver. Staging through two of our own pinned buffers instead lets the host memcpy of chunk N+1 overlap the DMA of chunk N (async copies from pinned memory are truly asynchronous), which approaches max(memcpy, PCIe) rather than their sum. The buffers are lazy-allocated on the first large upload and reused; gpu_staging_free releases them (a loader should call it once the weights are up — 128 MB of page-locked memory is not free).

: ~ i __gpu_stage_b 0

@ gpu_staging_free → v

: ~ i __gpu_reg_base 0

A host range the caller page-locked with gpu_host_register: copies whose source lies inside it are direct DMA already, so the staged path (an extra pass through DRAM) must NOT intercept them.

: ~ i __gpu_reg_size 0

@ gpu_host_register * u p i bytes → b

Page-lock an existing host range (an mmap'd model file) so uploads from it become direct DMA at PCIe rate — no host memcpy at all. At most ONE range is tracked; returns F when registration fails (not a CUDA backend, read-only registration unsupported, out of lockable memory) and uploads simply keep their staged path.

@ gpu_host_unregister * u p → 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

: ~ i g_params_buf 0

── 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. The void argument layer is rebuilt on EVERY launch — a decode step in a language model issues a few hundred of them, so a malloc/free pair per launch is pure overhead on the hot path. Keep one buffer and grow it on demand; a launch is synchronous with respect to reading the argument values (cuLaunchKernel copies them, cpu_launch reads them inline), so a single buffer per process is safe.

: ~ i g_params_cap 0

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

@ 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 @ cuDeviceGetAttribute *u out i32 attrib i32 dev → i32

cuDeviceGetAttribute: attribute 75 = COMPUTE_CAPABILITY_MAJOR, 76 = MINOR, 15 = TOTAL_CONSTANT_MEMORY … we need the first two to key compiled kernels by architecture and to rank devices.

& cuda @ cuDeviceTotalMem_v2 *u bytes 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 @ cuMemHostAlloc *u pp i bytes i32 flags → i32

& cuda @ cuMemFreeHost *u p → i32

& cuda @ cuMemHostRegister_v2 *u p i bytes i32 flags → i32

& cuda @ cuMemHostUnregister *u p → i32

& cuda @ cuMemcpyHtoDAsync_v2 i dst *u src i n i stream → i32

& cuda @ cuStreamSynchronize i stream → i32

& cuda @ cuStreamCreate *u out i32 flags → i32

& cuda @ cuStreamBeginCapture_v2 i stream i32 mode → i32

& cuda @ cuStreamEndCapture i stream *u graph_out → i32

& cuda @ cuGraphInstantiateWithFlags *u exec_out i graph i flags → i32

& cuda @ cuGraphLaunch i exec i stream → i32

& cuda @ cuGraphExecDestroy i exec → i32

& cuda @ cuGraphDestroy i graph → 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

& cuda @ cuEventCreate *u phe i32 flags → i32

Events — the only honest way to time a kernel. A host clock around a launch measures the launch, not the work; an event pair recorded on the launch stream measures the GPU. cuEventElapsedTime writes a float (4 bytes) into its out slot, not a double.

& cuda @ cuEventRecord i hev i stream → i32

& cuda @ cuEventSynchronize i hev → i32

& cuda @ cuEventElapsedTime *u pms i hstart i hend → i32

& cuda @ cuEventDestroy_v2 i hev → 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 @ nvrtcGetCUBINSize i prog *u sz → i32

& nvrtc @ nvrtcGetCUBIN i prog *u cubin → 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_device_cc i dev → i

Compute capability of a CUdevice as major*10 + minor (e.g. 89 for an RTX 4090's sm_89). 0 when the query fails.

@ cuda_device_mem i dev → i

Total device memory in bytes (0 on failure).

@ 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_compile_cubin s src s name i cc * u szout → *u

Compile straight to a device-specific CUBIN (-arch=sm_<cc>), so cuModuleLoadData needs no driver JIT — the JIT of a dozen PTX modules is seconds of pure start-up latency, every run. Returns 0 on ANY failure (old NVRTC without nvrtcGetCUBIN, unknown arch, compile error) and the caller falls back to the PTX path, which reports errors properly. szout (8 bytes) receives the cubin length — cubin is binary, there is no NUL to measure it by.

@ 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_host_alloc i bytes → *u

Page-locked (pinned) host memory: DMA-able, so an HtoD from it is a straight PCIe transfer with no driver-internal staging copy, and an ASYNC HtoD from it actually overlaps with host work. 0 on failure.

@ cuda_host_free * u p → i

@ cuda_host_register * u p i bytes → i

Page-lock an EXISTING host range (e.g. an mmap'd model file) so HtoD copies out of it are direct DMA instead of driver-staged. Flag 8 = CU_MEMHOSTREGISTER_READ_ONLY — required for PROT_READ mappings, which is exactly what a model file is. p must be page-aligned (mmap is).

@ cuda_host_unregister * u p → i

@ cuda_htod_async i dptr * u host i bytes i stream → i

@ cuda_stream_sync i stream → i

: ~ i g_cuda_stream 0

── 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. The stream launches ride. 0 (the legacy default stream) outside graph capture; during capture every launch must go to the CAPTURED stream — cuda_graph_begin points this at one, cuda_graph_end restores 0.

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

@ 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.

@ cuda_event_create → i

@ cuda_event_record i ev → i

@ cuda_event_sync i ev → i

@ cuda_event_elapsed_bits i start i end → i

Milliseconds between two recorded events, as the raw 32-bit float pattern the driver writes — the caller widens it (bits_to_f32). A float-typed out parameter is why this cannot just nurl_peek a double.

@ cuda_event_free i ev → v

: ~ i g_cuda_gstream 0

── CUDA Graphs: capture a launch sequence once, replay it as ONE call ── The per-node replay engine's cost on small graphs is pure launch overhead; a captured graph turns N launches into one cuGraphLaunch with the SAME kernels, argument values and order — bit-identical results, driver-level dispatch. The capture stream is created once and kept for the process (like the context); capture mode is THREAD_LOCAL so other threads' work cannot invalidate a capture.

@ cuda_graph_begin → i

Begin capturing: every cuda_launch until cuda_graph_end records into the graph instead of executing. 0 = ok.

@ cuda_graph_end → i

End the capture and instantiate an executable graph. Returns the exec handle (0 on failure). The interior CUgraph is destroyed after instantiation — the exec carries everything needed to launch.

@ cuda_graph_launch i exec → i

Launch the captured sequence and wait for it. 0 = ok.

@ cuda_graph_free i exec → v