Form-native GPU carrier lane for NVIDIA/RTX — CUDA twin of the Metal audit lane

NameDescription

The Form-native GPU codegen already exists: form-kernel-ts/src/backends/cuda.ts emits __global__ CUDA C++ (FP16 __hadd/__hmul, FP8 mma.m16n8k16 tensor-core, grid/block dispatch) and form-stdlib/jit-tensor-emit.fk emits the same kernels as Metal MSL (jte-matvec-msl, jte-affine-train-msl, mlp). The Metal lane is fully CLOSED end-to-end by scripts/metal_*_audit.sh: Form emits the kernel, swiftc+Metal.framework compiles & launches on the GPU, and a BIT-EXACT CPU parity gate (one thread per row, fp32 accumulation, mathMode=.safe) decides pass/fail. GAP on this Windows+RTX host: no CUDA host carrier exists. To run learning bands on the RTX we must (1) add CUDA emit recipes as twins of the MSL ones in jit-tensor-emit.fk (jte-matvec-cuda / jte-affine-train-cuda), or drive cuda.ts; (2) add scripts/cuda_matvec_audit.sh + cuda_affine_train_audit.sh — the dumb carrier: nvcc compile (CRITICAL: --fmad=false to forbid FMA contraction so mul-then-add stays two roundings and matches the CPU right-fold bit-for-bit), CUDA-driver/runtime host launcher (alloc device buffers, one-thread-per-row launch, copy back), same bit-exact parity gate, SKIP(2) when no nvcc/device; (3) add a windows-x64-cuda / sm_89 target row to form-stdlib/hati-os-targets.fk; (4) prereq: install CUDA Toolkit (nvidia-smi driver already present, nvcc is not). Proof model: GPU float is a CARRIER gated against the CPU oracle, not a four-way bit-identical band. SECOND AXIS (separate, cheaper): Ollama already runs on the RTX via its own ggml/CUDA backend (runner shells `ollama run`, no flags) — integrate it into the form-native model lane at the presence-host cmd=/kernel-walk boundary (native-first-else-walk), LLM output treated as a carrier checked against the recipe oracle; you do not lower Ollama through Form. Open questions: can bit-exact matvec parity hold on CUDA with --fmad=false + fp32 accum, or do we drop to an epsilon gate for backprop? Lift candidates from Modular Mojo open-source CUDA stack (Apache-2.0 parts only) for the host-launch carrier + numerics parity discipline — research in flight.

Potential value (USD)Estimated cost (USD)Confidence (0–1)