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: the opaque i64 device handles are plain fields (read them freely — . b dptr costs nothing), and each carries an own handle on the resource behind it. Nothing is released by hand: every copy of a Gpu / GpuKernel / GpuBuffer / GpuHost / GpuTimer / GpuGraph (a struct field, a Vec element, a capture) is the same resource, and its LAST owner releases it exactly as the explicit free did — device memory freed, the module unloaded, the event or graph destroyed, the context's retain released. Every resource also holds the context it lives in, so a context outlives its memory whichever order the owners go in. gpu_close / gpu_free / gpu_kernel_free / gpu_host_free / gpu_timer_free / gpu_graph_free remain as optional early releases of one owner. 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

: s GPU_WGSL_SENTINEL __nurl_wgsl_set``

The entry gpu_open probes to tell a working WebGPU host from a missing one. Every WGSL kernel set (web/kernels_wgsl.js) carries it, a no-op with no parameters, so the probe names no real kernel — those belong to whatever executor the set serves and change with it (it named osigmoid, a kernel of onnx's pre-0.7 set: a current set without it would read as "no WebGPU").

@ gpu_force_webgpu → v

: GpuCtxImpl { i dev i backend String name }

── ownership ───────────────────────────────────────────────────── A context retain taken by gpu_open, in an rcbox (stdlib/core/rcbox.nu): its last owner — the Gpu and everything allocated, compiled or created through it — gives back the pinned staging pair (it lives in this context) and releases the retain, which is what gpu_close did by hand. name caches the device name gpu_name hands out.

% Drop GpuCtxImpl { @ drop GpuCtxImpl x → v { ( __gpu_ctx_release . x dev . x backend ) } }

: GpuCtx { s ctl }

@ GpuCtx_drop sink GpuCtx h → v

: i GPU_RES_MEM 1 // device memory (cuMemFree / host RAM / GPUBuffer)

One device-side resource: h is the raw handle and kind says how to give it back, on the backend that made it. Holding the context keeps the context alive for as long as the resource is.

: i GPU_RES_MODULE 2 // a loaded module (cuModuleUnload)

: i GPU_RES_EVENT 3 // a CUDA event (cuEventDestroy)

: i GPU_RES_GRAPH 4 // an executable graph (cuGraphExecDestroy)

: i GPU_RES_HOST 5 // host staging memory (nurl_free)

: GpuResImpl { i kind i h i backend GpuCtx ctx }

% Drop GpuResImpl { @ drop GpuResImpl x → v { ( __gpu_res_release . x kind . x h . x backend ) } }

: GpuRes { s ctl }

@ GpuRes_drop sink GpuRes h → v

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

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

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

: GpuBuffer { i dptr i bytes GpuRes own } // 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_ok Gpu g → b

@ gpu_none → Gpu

A device that is not there (gpu_ok F) — opens nothing, probes no driver. For a holder that must carry a Gpu but never launches on it.

@ gpu_close sink Gpu g → v

Let go of g now rather than at the end of its owner's scope. The context goes when its last owner does — this Gpu, its copies, and every buffer, kernel, timer and graph made through it — and with it the pinned staging pair: that was allocated in this context, and a program that opens the device again (a server that unloads its model when idle) would otherwise hand the next upload two host buffers the driver no longer knows — a segfault inside the first cuMemcpyHtoDAsync, found exactly that way.

@ 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_bind_thread Gpu g → b

Make this device usable from the CALLING thread. On CUDA the context is thread-local, so a program that opens the device on one thread and launches from another has to say so; the CPU/static/WebGPU backends have no such state and answer T. Idempotent and cheap — call it at the top of any work item that may run on a thread the device was not opened on.

: GpuTimer { i ev GpuRes own }

── Timers (CUDA backend only; every other backend reports 0) ───────── A timer is one event recorded on the launch stream; a begin/end pair of them brackets the work between. Timing a launch from the host measures the launch; these measure the GPU. ev is 0 when the backend has no events, and every call below is a no-op on a 0 timer, so a caller need not branch on the backend. The event is destroyed with the timer's last owner.

@ gpu_timer_new Gpu g → GpuTimer

@ gpu_timer_none → GpuTimer

A timer that is not there — the placeholder for "not made yet".

@ gpu_timer_ok GpuTimer t → b

@ gpu_timer_mark Gpu g GpuTimer t → v

@ gpu_timer_ns Gpu g GpuTimer start GpuTimer 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 sink GpuTimer t → v

Let go of t now rather than at the end of its owner's scope.

: GpuGraph { i exec GpuRes own }

── 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. The executable graph is destroyed with its last owner.

@ gpu_graph_begin Gpu g → b

@ gpu_graph_end Gpu g → GpuGraph

The executable graph; not ok (gpu_graph_ok F) when capture or instantiation failed, or on a backend without graphs.

@ gpu_graph_none → GpuGraph

A graph that is not there — the placeholder for "not captured yet".

@ gpu_graph_ok GpuGraph x → b

@ gpu_graph_launch Gpu g GpuGraph x → i

@ gpu_graph_free sink GpuGraph x → v

Let go of x now rather than at the end of its owner's scope.

@ gpu_kernel_ok GpuKernel k → b

@ gpu_kernel_none → GpuKernel

A kernel that is not there (gpu_kernel_ok F): what a failed compile returns, and a placeholder for "not compiled yet".

@ gpu_kernel_free sink GpuKernel k → v

Let go of k now rather than at the end of its owner's scope. A CPU backend module is never dlclose'd (see cpu_module_free); a CUDA one is unloaded with the kernel's last owner.

: GpuHost { * u ptr i bytes GpuRes own }

── host (pinned-free) staging buffers ──────────────────────────── Plain host memory to stage data for upload / receive on download, released with its last owner. f32 is the GPU-native element type; these address it at 4-byte stride. ptr is the memory itself, for gpu_upload / gpu_download and anything that wants the raw address — valid while the GpuHost is.

@ gpu_host_none → GpuHost

No buffer — the placeholder for "not allocated yet".

@ gpu_host_ptr GpuHost h → *u

@ gpu_host_bytes GpuHost h → i

@ gpu_host_free sink GpuHost h → v

Let go of h now rather than at the end of its owner's scope.

@ gpu_buffer_view i dptr i bytes → GpuBuffer

A view of bytes of device memory at dptr that owns nothing — a sub-range of a live buffer, or memory something else allocated. The caller keeps the memory alive for as long as the view is used.

@ gpu_free sink GpuBuffer b → v

Let go of b now rather than at the end of its owner's scope.

: ~ 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

: GpuCopy { i dptr i host i bytes }

── one streamed upload for many tensors ───────────────────────────

A model is a thousand tensors and most of them are a few megabytes: below the staged path's 64 MB chunk, so each went up as its own synchronous copy out of pageable memory — the driver staging it through its own small pinned buffer, one thread, ~4 GB/s, nothing overlapping anything. gpu_upload_batch takes the whole list as ONE byte stream: it packs tensors back to back into the pinned staging pair, four threads filling a chunk while the DMA of the previous one drains (each buffer waits on its OWN event, so a fill never waits for the other buffer's copy), and issues one async copy per tensor out of the chunk. large-v3's 3.1 GB, 1,200 tensors: 0.65 s becomes ~0.3 s, bounded by PCIe rather than the driver.

host is an address (# i of the pointer) so an entry is a plain value. Every dptr must be a live allocation of at least bytes. 0 == success; on the CPU/WebGPU backends the list is uploaded one entry at a time.

@ gpu_upload_batch ( Vec GpuCopy ) items → i

: ~ 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_arg_buffer GpuBuffer b → i

@ gpu_arg_i32 i v → i

@ gpu_arg_i64 i 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_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 i32 flags → *u

The C ABI's ints are i32: on x86-64 an i64 in their place happens to work, but a wasm32 module that declares dlopen(…, i64) is INVALID — the call's operand type must match wasi-libc's signature, so every wasm build that linked this file failed to compile in the browser.

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

& c @ dlclose *u handle → i32

& c @ system s cmd → i32

& 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_module_free sink * 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_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.

: s CPU_STATIC_SENTINEL __nurl_static_set``

The registry: a name→launcher table + the strong nurl_static_kernel. The table always carries one extra row, CPU_STATIC_SENTINEL, bound to a no-op: gpu_open probes for it to tell "a static set is linked" from the runtime's weak NULL stub without naming any real kernel — kernel names belong to whatever library generated the set and change with it.

@ cpu_static_registry ( Vec String ) entries → String

---

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 @ cuMemGetInfo_v2 *u free *u total → 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

: ~ i __cuda_init_done 0

cuInit exactly ONCE per process, with the result remembered.

The driver's own contract says cuInit is idempotent, and on a working CUDA install a second call is a no-op. It is not harmless everywhere: in a sanitized build, where the first cuInit fails (ASan's allocator makes the driver's pinned-memory setup return CUDA_ERROR_OUT_OF_MEMORY), the SECOND call blocks for exactly 180 seconds and leaves a 408-byte allocation behind — which is what a gk_open_best (device probe, then open) turned every sanitized run of a GPU program into. One call, one answer, handed to every later caller.

: ~ i __cuda_init_rc 0

: ~ 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.

: ~ 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.