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

: 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_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 → i

& c @ nurl_poke_i32 *u base i idx i 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.

---

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 *u pctx i32 flags i32 dev → i32

& cuda @ cuCtxDestroy *u ctx → i32

& cuda @ cuCtxSynchronize → i32

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

& cuda @ cuModuleUnload *u module → i32

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

& cuda @ cuMemAlloc *u dptr i bytesize → i32

& cuda @ cuMemFree i dptr → i32

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

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

& cuda @ cuLaunchKernel *u f i32 gx i32 gy i32 gz i32 bx i32 by i32 bz i32 sh *u stream *u params *u 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) ──────────────────────────────────

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

& nvrtc @ nvrtcGetPTXSize *u prog *u sz → i32

& nvrtc @ nvrtcGetPTX *u prog *u ptx → i32

& nvrtc @ nvrtcGetProgramLogSize *u prog *u sz → i32

& nvrtc @ nvrtcGetProgramLog *u 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 → i

& c @ nurl_poke_i32 *u base i idx i 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 (borrowed copy into a fresh buffer).

@ cuda_ctx_create i dev → i

Create a context on a CUdevice → CUcontext handle (0 on failure).

@ 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_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").