OXIDE-001 (0.68.3, fast-track): cuda-oxide pure-Rust #[kernel] → PTX — one GDN kernel ported and measured on lambda-labs, gx10, yoga and mini; every kernel an ont:Kernel with parity, timing and safety receipts per silicon under pv lint --gate shapes

Operator 2026-09-19 (Noah): "for 0.68.3 I need you to fast track exploration of cuda-oxide as we discussed and ensure: lambda-labs, gx10, yoga and mini are able to test it, add as a ticket and also ensure we are using the pv ontology with it to guide performance and safety and correctness."

Where it stands (measured 2026-09-19): EPIC #3062 (NVIDIA CUDA Rust: cuda-core / cutile / cuda-oxide) parked at 0.73.0 with every row unchecked; no crate on main depends on cuda-oxide; one contract exists (contracts/cuda-oxide-rope-parity-v1.yaml: RoPE port, bit-parity vs f64 CPU on GB10, ≤ 1.2× hand PTX). The eight new GDN kernels (#3517) are hand-written PTX. Floors recorded then: cuda-core needs driver ≥ R580 (lambda-labs now 580.119.02), cutile needs CUDA 13.1+ (gx10 on 13.3, lambda-labs not), gx10's ptxas sits under the GH-480 rewriter.

Policy (KERNEL-001, enforced by the ontology, not by review)

New kernels are authored as cuda-oxide first, with a CPU f64 reference, and a parity + timing receipt on every CUDA host before merge.
A hand-PTX variant may exist only under a ptx_exemption naming the receipt (oxide slower than the stated margin, or a named missing intrinsic). The shape refuses an exemption without a receipt.
Existing hand PTX migrates in measured order: the GDN set first (memory-bound — conv1d, gates, norms, delta rule, partial RoPE, L2 norm), the Q4K/DP4A/tensor-core GEMV family last or never, by receipt.

The pv ontology half (with the ontology session, ONT-4b2's `extract:code`)

Every #[kernel] function is an ont:Kernel entity (a Rust symbol, so extract:code sees it; a PTX string is an opaque blob it cannot bind). Properties: kernel:reference (the f64 CPU function), kernel:authoring ∈ {oxide, ptx}, kernel:exemptionReceipt.
Shapes: kernel-parity — every ont:Kernel has a parity receipt per required CUDA host (cos ≥ 0.9999 ∧ max|Δ| < 1e-3 vs f64, the existing F-OXIDE-ROPE-PARITY-001 form); kernel-timing — oxide_us / handptx_us ≤ 1.2 per host, or an exemption receipt; kernel-safety — the oxide kernel is #![forbid(unsafe_code)]-clean at the Rust level, bounds-checked indexing, and its PTX passes ptxas with a recorded register budget (the O2 register_budget obligation of EPIC: nvidia-cuda-rust-library-integration.md — stabilize GPU quality with NVIDIA's CUDA Rust libraries (0.67) #3062, finally armed). Reported, then armed per shape (the ONT per-shape arming ratchet), never red-by-construction on main.
Receipts under evidence/kernels/<kernel>/<host>.json (schema apr-kernel-receipt/v1: kernel, host, cc, sha, parity, timing, register_budget, ptxas_version, authoring).

Four hosts, four lanes

host	silicon	what it tests
lambda-labs	sm_89, driver 580.119.02	parity + timing + ptxas budget (never a CI runner — receipts produced by the T-2 dogfood, committed)
gx10	sm_121, CUDA 13.3	same, on Blackwell; the GH-480 ptxas rewriter interaction recorded
yoga	RTX 4060 Laptop sm_89, driver 595	same on a second sm_89 with a NEWER driver: the driver-floor probe; runs on `yoga-eph` in CI
mini	M4, no CUDA	the compile half: `#[kernel]` → PTX codegen builds, PTX text golden tests, `ptxas`-free structural checks, `pv lint --gate shapes` over the kernel contracts; runs in the `mac-check` lane (#3488)
Each lane's first run is the first-green proof before its shape may block.

Rows

O-0 floors measured on all four hosts (driver, toolkit, rustc, ptxas) and recorded in this issue — no assumptions
O-1 pilot: gated_rmsnorm ported to cuda-oxide beside the hand PTX; parity + timing receipts on lambda-labs, gx10, yoga; PTX golden on mini
O-2 ont:Kernel entity + the three shapes in contracts/ontology.yaml / a kernel-*-v1 contract, reported first, armed per the ratchet (with the ontology session)
O-3 register_budget armed via load_sym! with a revert-to-RED proof (EPIC: nvidia-cuda-rust-library-integration.md — stabilize GPU quality with NVIDIA's CUDA Rust libraries (0.67) #3062 O2)
O-4 decision receipt: per GDN kernel, oxide or ptx_exemption, written from the receipts; EPIC: nvidia-cuda-rust-library-integration.md — stabilize GPU quality with NVIDIA's CUDA Rust libraries (0.67) #3062 re-milestoned here and closed or split
O-5 T-2 dogfood row: kernel receipts regenerated on the pinned binary before the bump

Refs #3062, #3519 (EPIC 0.68.3), #3517, contracts/cuda-oxide-rope-parity-v1.yaml, PMAT-921.

O-0 floors (read-only, all four hosts)

# O-0 floors, measured 2026-09-19T09:15:06Z
== lambda-labs (local)
host=noah-Lambda-Vector arch=x86_64 os=Linux
driver=580.119.02, NVIDIA GeForce RTX 4090, 8.9
toolkit=/usr/local/cuda-13.3 release 13.3 ptxas=V13.3.73
rustc=rustc 1.93.0 (254b59607 2026-01-19) nightly=3
llvm-nvptx=1 sccache=
cargo-registry-cuda-oxide=0
== gx10
host=gx10-a5b5 arch=aarch64 os=Linux
driver=590.48.01, NVIDIA GB10, 12.1
toolkit=/usr/local/cuda release 13.3 ptxas=V13.3.73
toolkit=/usr/local/cuda-13.3 release 13.3 ptxas=V13.3.73
toolkit=/usr/local/cuda-13.0 release 13.0 ptxas=V13.0.88
rustc=rustc 1.95.0 (59807616e 2026-04-14) nightly=2
llvm-nvptx=1 sccache=
cargo-registry-cuda-oxide=0
== yoga
host=yoga arch=x86_64 os=Linux
driver=595.91.07, NVIDIA GeForce RTX 4060 Laptop GPU, 8.9
toolkit=/usr/local/cuda release 13.3 ptxas=V13.3.73
toolkit=/usr/local/cuda-13.3 release 13.3 ptxas=V13.3.73
rustc=rustc 1.98.1 (48a229cea 2026-09-01) nightly=0
llvm-nvptx=1 sccache=
cargo-registry-cuda-oxide=0
== mini
host=mini arch=arm64 os=Darwin
driver=
rustc=rustc 1.98.1 (48a229cea 2026-09-01) nightly=1
llvm-nvptx=1 sccache=
cargo-registry-cuda-oxide=0

Notes: llvm-nvptx = whether the installed rustc lists the nvptx64-nvidia-cuda target (cuda-oxide emits PTX through it); the mini has no CUDA driver by construction and is the compile-half lane; cargo-registry-cuda-oxide = whether any crate source for cuda-oxide has ever been fetched on that host (0 everywhere = nothing has been built yet).

O-0 floors measured 2026-09-19 (all four hosts, `cargo oxide doctor` + a real device run)

cuda-oxide requirements per NVlabs README @ b9847e95 (2026-09-18): nightly-2026-08-28 + rust-src/rustc-dev/llvm-tools; CUDA toolkit 13.0+ (+ cuRAND headers); driver R580+ (libcuda loaded at run time); clang + resource headers for the host cuda-bindings bindgen (LLVM 21 only needed for TMA/tcgen05 lowering); Linux only (book §installation).

host	driver / GPU	toolkit used	ptxas	llc (rustup)	clang	nightly-2026-08-28 + 3 comps	cargo-oxide	`cargo oxide doctor`	`cargo oxide run` (1024-elem kernel)
lambda-labs	580.119.02 / RTX 4090 sm_89	13.3 (`/usr/local/cuda-13.3/bin` on PATH; the `cuda` symlink → 13.0 which has no nvcc, and bare `nvcc` resolves to apt's 12.8)	13.3.73	LLVM 23.1.0-rust	14.0.0 (accepted)	✓ (installed today)	b9847e95 (installed today)	✅ (nvJitLink reports 13.0 — mixed 13.0/13.3 libs)	PASSED
gx10	590.48.01 / GB10 sm_121	13.3	13.3.73	LLVM 23.1.0-rust	21.1.8	✓ (installed today)	was 43a76378 (June, 720 commits behind) → b9847e95 today	✅ after `update-alternatives clang → clang-21` (doctor probes bare `clang`; only `clang-21` existed)	PASSED
yoga	595.91.07 / RTX 4060 Laptop sm_89	13.3	13.3.73	LLVM 23.1.0-rust	21.1.8	✓ (installed today)	b9847e95 (installed today)	✅	PASSED
mini	none (M4, Darwin)	none	none	none	Apple clang 21	✓ (installed today)	not installed — Linux only upstream	n/a	n/a

Verdict: the runtime floor (driver ≥ R580, toolkit ≥ 13.0) is met on all three CUDA hosts, and each one compiles a #[kernel] to PTX and runs it on device with the current upstream head. Nothing was at the correct floor before today: no host had the nightly components, two had no cargo-oxide, gx10's was three months stale.

Findings to carry (not fixed here, they are declared-stack changes):

lambda-labs toolkit is split: /usr/local/cuda → 13.0 (no nvcc), apt nvcc 12.8 on PATH, nvJitLink 13.0 while nvcc/ptxas are 13.3. Needs a forjar row: update-alternatives --set cuda /usr/local/cuda-13.3 (attended, lambda-labs manifest) and the 12.8 apt nvcc removed or shadowed. Until then every oxide invocation on lambda must put /usr/local/cuda-13.3/bin first.
lambda-labs has no LLVM/clang 21 — fine for the GDN kernel class (no TMA/tcgen05), blocks nothing today; forjar row if Blackwell-class intrinsics are ever built here.
gx10 clang alternative was set ad hoc today (/usr/bin/clang → clang-21); needs the matching forjar symlink resource in machines/gx10 so it survives a converge.
mini's row must be re-scoped: upstream is Linux-only, so the mac-check lane can only do pv lint --gate shapes over the kernel contracts and PTX-text golden diffs of committed artifacts, not #[kernel] codegen. The nightly + components are installed there anyway.
Upstream trap for external projects: cargo oxide setup fails with no library targets found in package after it has already built and cached the backend; cargo oxide run alone is the correct path (it builds the backend from the project's git dependency on first use).
Pin: every receipt must record the cuda-oxide commit (b9847e95 today); upstream moves ~10 commits/day and gx10 silently sat 720 behind.

Milestone 0.70.0 → 0.70.1 (gate triage). 0.70.0 now holds only the operator's release gate: "FULL ontology spec and all aprender binary working" + "0.70 to crates.io" (ONT-001, every [[bin]] build+smoke, the publish cascade), plus P0s and operator-ruled rows. This issue isn't needed for any of those. Evidence: evidence/m070-gate-triage-1835.md. If that's wrong, move it back to 0.70.0 and say which gate it serves.

evidence(oxide): #3522 O-1 gated RMSNorm receipts — GO on lambda, yog…

fix(oxide): #3522 a timing NO-GO exited 0 — the gate was computed and…

evidence(oxide): #3522 lambda receipt at 5b29ed2 + timing-gate mutati…

O-1 pilot (gated_rmsnorm → cuda-oxide): GO on lambda, yoga, gx10. Branch feat/3522-oxide-gated-rmsnorm @1952808e3 (for fold; no PR).

Safe-Rust #[kernel] (no unsafe, DisjointSlice<LinearTiles<4>> writes), exp + ex2 variants, A/B against the committed hand PTX.
The hand PTX is pinned to the emitter by gdn_gated_rmsnorm_ptx_golden (runs with no GPU; this is the mini PTX-golden half).

host	cc	ratio exp/ex2	regs oxide/hand
lambda RTX 4090	sm_89	0.909 / 0.913	31 / 27
yoga RTX 4060L	sm_89	0.944 / 0.948	30 / 27
gx10 GB10	sm_121	1.000 / 0.998	32 / 29

Parity: cos 1.0000000, max|Δ| 1.43e-6 vs f64 on every host.
Receipts (apr-kernel-receipt/v1) are in evidence/kernels/gdn_gated_rmsnorm/; details in experiments/cuda-oxide/gated-rmsnorm/RESULTS.md.
Caveat: at this shape, both kernels sit on the launch floor, so the ratio shows no added cost; it cannot resolve arithmetic differences.
A first gx10 run read NO-GO (1.50). A foreign apr process was holding the GPU without the lock, and the hand PTX moved too. The clean re-run is the committed receipt.
Quorum: sonnet + haiku, degraded same-family. Round 1 found the timing gate never failed the exit status; this is fixed, and a mutation run confirms a NO-GO now exits 4. Round 2 was 2/2 APPROVE.

Next: O-2 (ont:Kernel + the kernel-parity / kernel-timing / kernel-safety shapes).

feat(ont): O-2 — ont:Kernel + kernel-parity/timing/ptx-exemption/safe…

fix(ont): kernel-safety missed unsafe METHODS, trait fns, extern bloc…

Milestone 0.71.0 -> 0.73.0 to match its epic E6 (train 0.73), rule 3. Dependency check (contracts on origin/main, pv status on contracts/ont-kernel-receipts-v1.yaml): the only references are contracts/binding.yaml (a notes line) and ont-kernel-receipts-v1 (F-KRC-001), whose kernel-parity/timing/safety shapes are REPORTED and not armed; arming follows this ticket's first receipt and is a new GATE only by C324 rule 7. Nothing in the 0.71.0 floor depends on it. Labels already say parked, owner la-0.73. Agent: aprender-triage-worker

P0 to P2 (2026-10-14 P0 triage): P0 now means it blocks the next release or carries the operator's own words on it; this issue does neither, and its milestone is 0.73.0.

merged into #5313 (master ticket for epic #3999; triage rule 7b, 2026-10-14). It is a checklist row there; reopen this issue when work on it starts or its PR opens.

Moved: this sub-ticket is now a row in #5338 (Backend portability: correct kernels and config on every CPU, wgpu and CUDA target), split out of #5313 by theme.

Navigation Menu

Repository navigation

OXIDE-001 (0.68.3, fast-track): cuda-oxide pure-Rust #[kernel] → PTX — one GDN kernel ported and measured on lambda-labs, gx10, yoga and mini; every kernel an ont:Kernel with parity, timing and safety receipts per silicon under pv lint --gate shapes #3522

Policy (KERNEL-001, enforced by the ontology, not by review)

The pv ontology half (with the ontology session, ONT-4b2's `extract:code`)

Four hosts, four lanes

Rows

O-0 floors (read-only, all four hosts)

O-0 floors measured 2026-09-19 (all four hosts, `cargo oxide doctor` + a real device run)

47 remaining items

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

OXIDE-001 (0.68.3, fast-track): cuda-oxide pure-Rust #[kernel] → PTX — one GDN kernel ported and measured on lambda-labs, gx10, yoga and mini; every kernel an ont:Kernel with parity, timing and safety receipts per silicon under pv lint --gate shapes #3522

Description

Policy (KERNEL-001, enforced by the ontology, not by review)

The pv ontology half (with the ontology session, ONT-4b2's extract:code)

Four hosts, four lanes

Rows

Activity

noahgift commented on Sep 19, 2026

O-0 floors (read-only, all four hosts)

noahgift commented on Sep 19, 2026

O-0 floors measured 2026-09-19 (all four hosts, cargo oxide doctor + a real device run)

noahgift commented on Sep 24, 2026

noahgift commented on Sep 25, 2026

47 remaining items

noahgift commented on Oct 9, 2026

noahgift commented on Oct 14, 2026

noahgift commented on Oct 14, 2026

noahgift commented on Oct 14, 2026

Metadata

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

Issue actions

The pv ontology half (with the ontology session, ONT-4b2's `extract:code`)

O-0 floors measured 2026-09-19 (all four hosts, `cargo oxide doctor` + a real device run)