[dependencies]
gpu = "^0.17.7"

gpu — GPU compute for NURL

A backend-neutral GPU compute interface for NURL, with four backends: CUDA, CPU (runtime-compiled), static-kernel, and WebGPU. You write a kernel in CUDA-C, and it runs on the GPU (NVRTC compiles it to PTX at runtime; the CUDA Driver API launches it — bound directly from pure NURL, no C bridge in the core runtime), on the CPU when no GPU is available, from a set of precompiled kernels linked into a sealed binary, or — in a browser/Deno wasm module — as WGSL compute shaders on the browser GPU (backend 3).

WebGPU backend (backend 3) — the browser GPU

( gpu_force_webgpu ) before gpu_open selects a WebGPU backend: a fixed set of WGSL kernels (web/kernels_wgsl.js) run through navigator.gpu driven by host imports the JS embedder implements (web/webgpu.js, works in a browser worker and in Deno). Device memory is a GPUBuffer; the one async op — GPUBuffer.mapAsync readback — is bridged with Asyncify so the synchronous NURL gpu_download returns data in wasm memory (web/wgpu_asyncify.c provides the unwind stack; build with wasmbuilder --asyncify-imports env.wgpu_download). A buffer's handle — the device pointer a program holds — is its id shifted up 32 bits, so a view (pointer plus a byte offset, as on every other backend) names the same buffer at that offset; each kernel argument reads and writes from its own.

Kernels are looked up by entry name — a checked gpukit kernel by its base name, the entry name without its __gk<hash> suffix — so the set must be exactly what the executor requests. For the onnx executor it is checked against the executor itself: each entry states the C parameter list it implements (sig) and the WGSL name of each buffer (names), and the bindings, uniform block and argument decoding are generated from that list; onnx's tests/wgsl_census_test.nu (no device) fails when the executor's kernel census holds a kernel the set lacks, a sig that differs from the recorded parameter list, buffer names that are not the bindings the kernel reads, two kernels under one base name, a parameter type the glue cannot marshal, or an entry the census no longer records. Only the WGSL bodies are written by hand. gpu_open probes the set's sentinel entry (GPU_WGSL_SENTINEL), never a real kernel name.

Every WGSL kernel runs on a real WebGPU device through this host, against a JS reference of its CUDA-C (./tests/webgpu_test.sh — Deno, or headless Chrome via tests/webgpu_chrome.mjs, on the GPU or SwiftShader; skips without either). onnx's tests/webgpu_test.sh runs a whole model forward in wasm on it. This is what runs the YOLOE detector in the browser (see packages/yoloe-demo).

CPU backend (v0.2.0) — run with no GPU

gpu_open selects CUDA when a device is present, and otherwise falls back to the CPU backend (NURL_GPU=cpu forces it). The CPU backend runs the same CUDA-C kernels: each is wrapped in a small CUDA-compatibility shim (blockIdx / threadIdx / blockDim / gridDim as thread-locals, __global__ / __device__ 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 neutral Gpu / GpuKernel / GpuBuffer surface is unchanged and every caller (onnx, objdet, yoloe) gets the fallback for free. The only host requirement is a C++ compiler on PATH — the CPU analogue of NVRTC.

Verified identical results CPU vs GPU: the onnx tiny MLP matches the onnxruntime reference, and the full 267-node YOLOE-seg forward gives the same detections on the CPU as on an RTX 4090.

On a machine with no NVIDIA driver at all, the toolchain links small stub objects (stdlib/{cuda,nvrtc}_stubs.o) in place of -lcuda / -lnvrtc, so a packages/gpu program still links and loads and auto-falls-back to the CPU. Build a portable CPU-only binary even on a GPU host with NURL_GPU_STUBS=1.

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

CUDA devices: 2
Using device 0: NVIDIA GeForce RTX 4090

[vadd] c[i] = a[i] + b[i]
  OK — 1000000 elems, spot c[7]=21 (=3*7)

[saxpy] y[i] = alpha*x[i] + y[i],  alpha=2.5
  OK — verified vs CPU

All GPU kernels verified. ✓

Design — why it stays out of the core

NURL's core (runtime.c) installs unchanged on machines with no GPU (e.g. a MILK-V Duo). So this package touches none of it:

• The CUDA Driver API (& \cuda\`) and NVRTC (& \nvrtc\`) are bound

with inline FFI declarations in src/cuda.nu — the only file that imports anything external.

• The toolchain auto-links -lcuda / -lnvrtc **only when a program

actually references those symbols** (the same mechanism used for opus / ALSA / SDL2). A GPU-less program never grows the dependency.

• libcuda ships with the NVIDIA driver — no CUDA toolkit needed at

runtime. Runtime kernel compilation needs NVRTC (libnvrtc, from the toolkit); a program that embeds pre-built PTX needs only the driver.

The only core addition is four generic, dependency-free typed-buffer accessors in runtime.c (nurl_peek_f32 / nurl_poke_f32 / _i32) — 4-byte-stride access to the float32 / int32 arrays that are the native GPU element layout (the stock 8-byte nurl_peek/poke can't address them). They ship to every target and require no GPU.

API (src/gpu.nu)

Backend-neutral types — Gpu, GpuKernel, GpuBuffer — so a future backend (ROCm/HIP, OpenCL, a CPU fallback) slots in behind the same names.

function	purpose
gpu_device_count → i	number of CUDA devices
gpu_open i ordinal → Gpu	init driver, bind device, retain the PRIMARY context (shared process-wide)
gpu_ok Gpu → b / gpu_name Gpu → s	status / device name
gpu_compile Gpu s src s name → GpuKernel	NVRTC compile + load extern "C" __global__ entry (raw)
gpu_alloc Gpu i bytes → GpuBuffer	device allocation
gpu_host_alloc i bytes → GpuHost	zero-filled host staging buffer (+ bounds-checked gpu_host_{set,get}_{f32,i32}; gpu_host_ptr for the raw address, raw)
gpu_host_f32 GpuHost → (Slice f32) / gpu_host_u8 → (Slice u)	the buffer as a checked view that lives no longer than it — the fast safe path
gpu_host_of_vec sink (Vec u) → GpuHost	adopt a byte Vec as a host buffer, no copy
gpu_upload_host GpuBuffer GpuHost → i / gpu_download_host GpuHost GpuBuffer → i	whole-buffer copies, safe: a host side of the wrong size is refused (-1)
gpu_upload_vec GpuBuffer (Vec u) → i / gpu_download_vec GpuBuffer → (Vec u)	the same for a byte Vec
gpu_upload GpuBuffer *u host → i	host → device from a raw address (the caller vouches it holds the buffer's size)
gpu_upload_batch ( Vec GpuCopy ) items → i	many host spans → many device buffers in ONE streamed pass through the pinned staging pair (a model's thousand tensors: PCIe-bound, not driver-bound); GpuCopy { i dptr i host i bytes }, host is the pointer as # i (raw)
gpu_download *u host GpuBuffer → i	device → host
gpu_arg_buffer/_i32/_i64/_f32 → i	encode one kernel argument
gpu_launch GpuKernel i grid i block (Vec i) args → i	1-D launch (raw)
gpu_grid i n i block → i	ceil-div grid size
gpu_sync Gpu → i	block until work completes
gpu_bind_thread Gpu → b	make the device current on THIS thread (CUDA contexts are thread-local; no-op T elsewhere)
gpu_buffer_view i dptr i bytes → GpuBuffer	an unowned view of device memory something else keeps alive (raw)
gpu_timer_new / gpu_graph_end → GpuTimer / GpuGraph	CUDA events / executable graphs
gpu_free / gpu_kernel_free / gpu_host_free / gpu_timer_free / gpu_graph_free / gpu_close	early release (optional)

What is raw (grammar §3.3e) trusts its caller for what the compiler cannot check — that kernel source is sound, that an address is live device memory of the size a kernel indexes — so only an unsafe function may call it: the one that compiled the kernel, allocated the buffers and checked the sizes. The handle structs are raw too: a GpuBuffer comes from gpu_alloc, never from a number. Everything else is safe to call, and tests/holes_test.sh holds the ways safe code used to reach memory it did not own.

Nothing is released by hand. Every handle (Gpu, GpuBuffer, GpuKernel, GpuHost, GpuTimer, GpuGraph) is a small value struct whose raw fields (. b dptr, . k func, …) read for free and whose own field holds the resource: every copy of it — a struct field, a Vec element — is the same resource, and its LAST owner releases it exactly as the explicit free did (device memory freed, module unloaded, event / graph destroyed, the context's retain released). Each resource also holds its context, so a buffer may outlive the Gpu it came from.

Example

$ `gpu.nu`

// unsafe: compiles one fixed kernel and launches it on buffers of n floats
unsafe @ main → i {
    : Gpu g ( gpu_open 0 )
    : GpuKernel k ( gpu_compile g
        `extern "C" __global__ void vadd(const float* a, const float* b, float* c, int n){
             int i = blockIdx.x*blockDim.x + threadIdx.x; if (i<n) c[i]=a[i]+b[i]; }`
        `vadd` )

    : i n 1024
    : i bytes * n 4
    : GpuHost ha ( gpu_host_alloc bytes )  // fill ha, hb through gpu_host_f32 ...
    : GpuBuffer da ( gpu_alloc g bytes )
    ( gpu_upload_host da ha )  // ... db, dc likewise

    : ( Vec i ) args ( vec_new [i] )
    ( vec_push [i] args ( gpu_arg_buffer da ) )   // a, b, c, then:
    ( vec_push [i] args ( gpu_arg_i32 n ) )
    ( gpu_launch k ( gpu_grid n 256 ) 256 args )
    ( gpu_sync g )
    ( gpu_download_host hc dc )
    ^ 0                                    // everything above goes with its owner
}

Kernel-argument ABI

gpu_launch takes a Vec i of arguments built with the gpu_arg_* encoders. Each encodes one value into an i64 cell; gpu_launch builds the void** array CUDA's cuLaunchKernel expects, pointing each entry at its cell. CUDA reads sizeof(param) bytes, so a 4-byte int/float in the low bytes of an 8-byte cell is correct on little-endian. A float scalar must go through gpu_arg_f32 (it stores the 32-bit IEEE-754 pattern); passing a raw double would mis-encode it.

Requirements

• NVIDIA driver (libcuda.so) — for any GPU work.
• NVRTC (libnvrtc.so, CUDA toolkit) — for gpu_compile. Embed PTX to

avoid it.

• Tested against CUDA 13 driver / CUDA 12 toolkit on GTX 970 + RTX 4090.

Status & roadmap

v0.1.0 is a PoC: 1-D launches, synchronous execution, f32/i32/pointer arguments. Not yet: streams/async, multi-dimensional grids beyond x, shared-memory configuration, a second backend. The interface is shaped to add them without breaking callers.