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

@ main → i

---

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 in *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

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

: GpuCtx { ( Arc GpuCtxImpl ) a }

@ GpuCtx_share GpuCtx h → GpuCtx

@ 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 (the resource's own Vec)

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

: GpuRes { ( Arc GpuResImpl ) a }

@ GpuRes_share GpuRes h → GpuRes

@ GpuRes_drop sink GpuRes h → v

@ 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_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_timer_ok GpuTimer t → b

@ gpu_timer_free sink GpuTimer t → v

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

@ gpu_graph_begin Gpu g → b

@ gpu_graph_ok GpuGraph x → b

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

: ( Locked i ) g_gpu_stage_lk ( __gpu_stage_lk0 )

@ gpu_staging_free → v

Give the pinned staging pair back (a loader calls this once its weights are up). Waits for any staged upload still running.

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

@ gpu_arg_i32 i v → i

@ gpu_arg_i64 i v → i

: ~ ( Vec i ) g_params 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.

: ~ b g_params_made F

@ 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 in *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

---

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 inout i32 count → i32

& cuda @ cuDeviceGet inout i32 device i32 ordinal → i32

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

& cuda @ cuDeviceGetAttribute inout i32 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 inout i bytes i32 dev → i32

& cuda @ cuCtxCreate_v2 inout i pctx i32 flags i32 dev → i32

& cuda @ cuDevicePrimaryCtxRetain inout i 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 inout i module *u image → i32

& cuda @ cuModuleUnload i module → i32

& cuda @ cuModuleGetFunction inout i hfunc i hmod s name → i32

& cuda @ cuMemAlloc_v2 inout i dptr i bytesize → i32

& cuda @ cuMemGetInfo_v2 inout i free inout i total → i32

& cuda @ cuMemFree_v2 i dptr → i32

& cuda @ cuMemcpyHtoD_v2 i dst in *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 inout i 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 in *u src i n i stream → i32

& cuda @ cuStreamSynchronize i stream → i32

& cuda @ cuStreamCreate inout i out i32 flags → i32

& cuda @ cuStreamBeginCapture_v2 i stream i32 mode → i32

& cuda @ cuStreamEndCapture i stream inout i graph_out → i32

& cuda @ cuGraphInstantiateWithFlags inout i 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 inout i pstr → i32

& cuda @ cuEventCreate inout i 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 inout f32 pms i hstart i hend → i32

& cuda @ cuEventDestroy_v2 i hev → i32

& nvrtc @ nvrtcCreateProgram inout i 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 inout i opts → i32

& nvrtc @ nvrtcGetPTXSize i prog inout i sz → i32

& nvrtc @ nvrtcGetPTX i prog *u ptx → i32

& nvrtc @ nvrtcGetCUBINSize i prog inout i sz → i32

& nvrtc @ nvrtcGetCUBIN i prog *u cubin → i32

& nvrtc @ nvrtcGetProgramLogSize i prog inout i sz → i32

& nvrtc @ nvrtcGetProgramLog i prog *u log → i32

& nvrtc @ nvrtcDestroyProgram inout i prog → i32

& c @ nurl_peek_f32 in *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 in *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.