minfer CUDA Backend Design

How minfer runs the compute graph on an NVIDIA GPU: the CudaBackend graph executor, the cuda.rs device layer, the kernel families it dispatches, the CUDA Graph capture/replay machine, and the memory and safety rules that hold it together.

Status. Landed. Phase 7 (7a–7e) shipped the backend, and the later campaigns (Phase 8, R/P5, the MMQ line, the decode campaign D1–D4, D5-R speculative decoding) built on it. Every mechanism described here is implemented in the tree. Baseline: HEAD = 916a7d2 (2026-09-14).

Provenance. This file was docs/CUDA-BACKEND-PLAN.md, written before Phase 7 as a plan. The section skeleton is preserved; the body has been rewritten in the present tense against the current code, and the plan-time "current state" / "placeholder" / v1-matrix text has been replaced by the landed design. Plan-era commit hashes that no longer resolve are noted where cited.

Related records. This document is the backend design (what it is and how it fits the graph). The optimization campaign — every measured lever, accepted or reverted — lives in docs/CUDA_OPTIMIZATION.md (live state + §0 history table + Appendix A env-gate reference) and its expansion layer docs/cuda_optimization_steps/ (one document per step). docs/CUDA-TECH-PRIMER.md explains the GPU techniques, docs/GPU_SAFETY.md holds the hard safety rules (CUDA section), and docs/inference_e2e_walkthrough/15-cuda-backend.md narrates the backend for a first-time reader. The graph contract this backend implements is docs/COMPUTE-GRAPH-DESIGN.md §3.5/§9.

1. Design Goals and Outcome

1.1 Goal

Implement CudaBackend (src/graph/cuda_backend.rs) as the third backend of the compute graph by wrapping the existing src/cuda.rs device layer — not stubbing it — so that on a CUDA machine the whole per-layer chain runs on the GPU through the standard build → assign → fuse → alloc → execute pipeline, with the same correctness contract as CPU and Metal: backend placement is decided at build time, kernel-invariant violations return Err, and there is never a silent mid-run fallback.

1.2 Outcome

1.3 Non-goals

• Multi-GPU / tensor split (llama.cpp's split machinery) — single-GPU engine.
• FP16/BF16 activations and cuBLAS/cublasLt. Prefill uses f16 weight tiles and int8 tensor-core MMQ, and the KV cache can be f16, but activations stay f32 and the matmul path is minfer's own.
• IQ*/Q2_K/Q3_K kernels — not implemented and not planned without a model that needs them.
• Windows and a self-hosted CUDA CI runner (deferred; device-gated tests skip gracefully).
• graph_optimize-style node reordering — the graph topology is the builder's decision.

1.4 Related records

2. Architecture at a Glance

2.1 Three layers

The split is deliberate: cuda.rs is the only place that touches the CUDA runtime API, and cuda_backend.rs is the only place that knows about graph nodes.

2.2 The CudaBackend surface

#![allow(unused)]
fn main() {
pub struct CudaBackend {
    state: &'static crate::cuda::CudaState,   // process-wide device context
    stream: *mut c_void,                      // THIS instance's non-blocking stream (#188)
    kv_layout: i32,                           // KV layout tag for this instance
    pool: Vec<CudaBuf>,                       // id -> { ptr, bytes }
    free: Vec<usize>,                         // byte-length-matched free list
    pool_gen: u64,                            // bumped on every pool allocation
    pos_scratch: *mut c_void,                 // device i32 positions plane
    pos_scratch_bytes: usize,
    pos_memo: Option<(usize, u64)>,           // one-execution-window conversion memo
    graph_execs: Vec<CapturedGraph>,          // instantiated graphs
    graph_runs: HashMap<(u64, (usize, usize)), u32>, // warmup counters
    capturing: Option<(u64, (usize, usize))>, // open capture window (on `stream`)
    graphs_mode: GraphMode,                   // Enabled / Disabled
    prefill_capture: bool,                    // default ON (MINFER_NO_PREFILL_CAPTURE opts out)
    cap: crate::cuda::CaptureStaging,         // viz/trace async D2H staging
}
}

new() returns None when CudaState::get() fails (no device, or MINFER_DISABLE_CUDA=1, both handled inside the device layer), which is how GraphAllocator::enable_cuda declines. The struct is unsafe impl Send/Sync: raw device pointers are dereferenced only by the GPU, and mutation happens only through &mut self (mirroring MetalBackend).

2.3 Weight residency and registry

Every graph-referenced tensor is uploaded once at model load and referenced by name afterwards:

• CudaState::register_weight(name, data) allocates and H2D-copies; the same name and size is a no-op (reloading the same GGUF must not leak a second device model), while a different size replaces the entry and deliberately leaks the stale buffer, because a live captured graph may still reference it. The leak is bounded by the number of distinct (architecture, tensor) shapes.
• has_weight_of_size(name, raw_len) is the size-aware gate used by the model wiring; it matches repacked weights by their original raw length, so a foreign-architecture entry reads as "not registered" and that model stays on CPU.
• get_weight_ptr(name) is the only lookup the executor uses.
• Quant-specific registrations: Q6_K uses a 224-byte-padded repack (register_weight_q6k_padded, enabling 16-byte uint4 loads); Q8_0 optionally registers a p32 split plane (register_weight_q80_p32); Q6_K/Q4_K optionally build pre-expanded / pre-decoded planes (register_weight_q6k_exp, register_weight_q6k_dsc, register_weight_q4k_dsc, and the D4-4 {name}__dpl dense split plane). Plane maps are keyed by the weight's device pointer and looked up inside prefill_mmq.
• The q4_K W_dsc plane is admitted for q4_K and only q4_K (src/q4k_dsc.rs, q4k_dsc_plane_admitted — issue #165). The rule has two halves: the type gate (TensorType::Q4_K is the one type mmq_raw_nb_bt dispatches the dsc template for, so the loader admits no other type) and the payload gate (raw.len() must be exactly od * (id / 256) * 144, q4_K's own block layout — equality, not a lower bound). register_weight_q4k_dsc re-checks the payload before the budget query and before expand_q4k_dsc, and expand_q4k_dsc itself returns None for a payload it cannot index, so a direct caller cannot bypass either. Why both: a q4_0 payload has exactly q4_K's bytes/element ratio (18/32 == 144/256), so the size check cannot refuse it; a q8_0 payload (34/32) is longer and would be misread as 144-byte q4_K super-blocks; and a future type with a smaller ratio (a 2-bit K-quant: 84/256) is shorter than the row arithmetic needs, so the size check is what refuses it instead of reading past the tensor. The check cannot tell a q4_K payload from another type's bytes of the same length — that is the type gate's job.
• The per-tensor registration dispatch is one shared rule (src/models/weight_reg.rs, issue #167). Both loaders call register_cuda_weight for every tensor the E5 plan puts on the device; it carries the whole contract: the quantized-type matches! set, the F16 raw branch, the F32 (1-D norms/biases vs 2-D matmul weights) branch, the Q6_K padded repack, the q8_0 p32 plane, the q4_K W_dsc plane under q4k_dsc_plane_admitted, and the clear_mmq_nb_bt_only rule. The decision (cuda_weight_reg) is pure — no CudaState, no environment; the r59 dispatch gates are passed in — so CI's CPU job runs its tests, exactly like src/q4k_dsc.rs. Both loaders previously carried a copy of this block and the copies had drifted twice: the qwen3 copy had neither the f16 branch (#141) nor the q4_K dsc call (r59/#165), so an f16 Qwen3 fell to the CPU and a q4_K Qwen3 kept the in-kernel scalar dsc decode. The graph-side type gate is per architecture and must list the same types (Qwen3Graph::weights_on_cuda gained F16 in #167).
• A per-weight f16 dequant cache (w16_cache) is enabled by the loader only when quantized matmul weights exceed 2 GiB and MMQ is off; MINFER_NO_W16CACHE=1 reverts.
• ModelLoadGuard (reentrant, process-wide) serializes loader registration so two models with same-named tensors cannot interleave, and real-model tests hold it across their forwards.

The graph-side CPU registry is separate: register_graph_weights registers into the allocator's CPU backend, which is what a CPU split of a mixed graph executes from; the CUDA registry is filled by the loaders at model-load time.

2.4 Streams, capture windows and the per-instance discipline

Everything the device path issues — kernels, cudaMemcpyAsync staging, events, capture/replay, synchronize — runs on a stream owned by the CudaBackend instance, not on a process-wide one. That is the #188 change:

• CudaState stays the process-wide context: the device, the name-keyed weight registry, the derived weight planes (q6k_exp/q6k_dsc/q4k_dsc/w16_cache) and the device_memory/host_alloc queries are genuinely context-scoped and shared.

• CudaState::create_stream() returns a fresh cudaStreamNonBlocking stream. CudaBackend::with_layout creates one per instance (no stream, no backend) and Drop destroys it after the pool, the scratches and the captured graphs.

• Every CudaBackend device operation starts with let _bound = self.bind(). bind_stream publishes the instance's stream in a thread-local; CudaState::stream() answers with it, so the ~60 launch/copy/event helpers in cuda.rs keep their signatures and still follow the instance. The stream consumers, named:

  Consumer	How it gets its stream
  kernel launchers (cuda.rs, the MMQ/attention/fused families)	self.stream() → the bound instance stream
  H2D input fill (write_input_async), D2D (copy_device_to_device), D2H staging (copy_to_host_async)	self.stream(); the pinned staging ring is keyed on the stream too
  events (record_event, stream_wait_event) and the F5 copy_cross/await_cross hooks	self.stream() via enqueue_cross_host/take_cross, both of which bind
  cudaGraphLaunch in graph_replay_step, graph_begin_capture, graph_end_capture_to_exec	self.stream() while the backend's own window is open
  synchronize / state_sync	self.stream() — one stream sync, counted per backend
  copy_cells (kv_move_rows) and alloc_buffer/free_buffer	the bound stream (the cudaFree/cudaMalloc themselves are context-wide)
  Drop for CudaBackend	binds, then frees pool/scratch/host/graph, then destroys its stream
  weight registration (CudaState::register_weight)	the context stream, never a backend's: an H2D copy queued on the context stream + cudaStreamSynchronize(context), so it is stream-ordered and cannot be recorded into anybody's capture window
  context-wide, stays shared	cudaMalloc/cudaFree, cudaMemGetInfo (device_memory), cudaHostAlloc/cudaFreeHost (host_alloc/host_free), the weight registry and the derived planes, cudaGetLastError

• Activation scratch is per stream. The buf_hidden/buf_q8_prefill/buf_qa8_t/ buf_q8_decode/buf_attn_partial/… slots are now StreamScratch, a map keyed on the current stream, and the MmqCache memo is keyed the same way (a hit records a scratch pointer, so a shared memo would hand one engine's plane to another). The staging ring (write_input_async) is keyed on the stream for the same reason. Everything unbound — the legacy layer path, direct CudaState tests — keys on the context stream, which is why the #185 guard is narrowed to that path (§2.5).

Capture mode. Windows are opened with cudaStreamCaptureModeThreadLocal (graph_begin_capture), not cudaStreamCaptureModeGlobal. Under Global another thread's capture-unsafe driver call belongs to the window: it either invalidates the capture (cudaErrorStreamCaptureInvalidated, 901) or faults inside the driver. The recorded SIGSEGV (#185) was exactly that — one thread at cuMemcpyHtoD_v2 under register_weight while another was at cuGraphInstantiateWithFlags under graph_end_capture_to_exec. Thread-local mode scopes invalidation to the capturing thread, so a registration on another thread is benign. MINFER_CUDA_CAPTURE_MODE=0|1|2 overrides the mode (relaxed / global / thread-local) — it exists for the #188 probe's measurement, not for production.

The #188 probe is the instrument. graph::cuda_backend::tests:: capture_window_on_one_thread_survives_a_weight_registration_on_another opens a capture window on one thread and issues a weight-registration copy on another while the window is open, then closes it and checks both that cudaStreamEndCapture returned 0 (never 901) and that the replayed graph produced the bytes it recorded. It has two env knobs so the mode can be judged rather than assumed: MINFER_PROBE_STREAM=context captures on the context (blocking) stream — the pre-#188 shared-stream model — and MINFER_PROBE_LEGACY_MEMCPY=1 issues the registration with the pre-#188 blocking cudaMemcpy. Measured on GB10 sm_121, 5 process runs per cell (90 s watchdog; crash = the probe's own assertion failed), 2026-09-27:

Three measured readings, none assumed:

1. The blocking copy is a hard deadlock, independent of the mode. A blocking cudaMemcpy is issued on the legacy null stream, which implicitly synchronizes with every blocking stream — including the one holding the open capture window, which by construction cannot complete until the host closes it. All three modes hang 5/5. So the mode is not the fix for the historical setup; a stream-ordered copy on a non-blocking instance stream is.
2. Relaxed is ruled out by direct measurement. With the structural fix in place, relaxed still returns cudaErrorStreamCaptureInvalidated (901) — the exact code the acceptance forbids — in 5/5 runs (cudaMalloc inside the window also fails, CUDA: failed to allocate 16384 bytes). It is not adopted.
3. Thread-local is the mode. With it, the probe passes 5/5 in both the pre-#188 shared-stream cell (the deadlock aside) and the instance cell; global also passes the instance cell but is the mode that lets a foreign thread's driver call belong to the capture, which is the class this ticket exists to remove.

Two readings: the mode is what makes the historical shared-stream setup safe (global is not viable; thread-local and relaxed both are, and thread-local keeps the capturing thread's own mistakes fatal, so it is the one adopted), and the structural change makes the mode irrelevant by removing the sharing. The concurrent device gate (models::qwen2::graph::tests::cuda_kv::two_cuda_engines_forward_concurrently_and_stay_bitwise_identical) is the positive half: two engines on two threads, two distinct streams, bitwise equal to their serial references.

The prefill-GEMM smem opt-in (#145, #147; lazy and gated since #218). A gemm_f16_nt_kernel_t instantiation whose dynamic shared memory exceeds the 48 KiB default must be opted in with cudaFuncSetAttribute(.., cudaFuncAttributeMaxDynamicSharedMemorySize, N) — a launch over an un-opted-in dynamic smem is rejected with cudaErrorInvalidValue and cannot succeed. The number requested is the kernel's own byte layout — `As 2TNKS halves + Am 2TNKS floats (AF32 only)

• Bs 2TMKS halves + Cs NW*256 floats, TN = 64, NW = blockDim.x/32= 8 — andgemm_dynamic_smem_bytes(tm, ks, af32) is the single source the launcher (launch_gemm_f16) reads; before #145 an eager sweep carried a stale copy of it while the launcher carried a copy that dropped the AF32 mirror. A request that exceeds the device's own cudaDevAttrMaxSharedMemoryPerBlockOptinis **skipped with the reason printed**, instead of called: the call could only returncudaErrorInvalidValue(whichcompute-sanitizercounts) and the instantiation cannot launch on that device at all. On GB10/sm_121 (limit 101376 B) that is exactly one combination,gemm_f16_nt_kernel_t<256,64,true>` at 122880 B.

The design is eager pre-warm at context creation + lazy per-launch opt-in (#223 restored the eager half). #188 deleted gemm_prefill_smem_init's CudaState::try_new call site with no mention in its commit message or its docs commit; two later "dead code hygiene" commits annotated the orphan #[cfg_attr(not(test), allow(dead_code))] instead of asking why a production-looking init had no production caller; #218 removed the function and its checked/skipped introspection, leaving the invariant tested but not enforced by production. #223 put the runtime guarantee back at the same site: CudaState::try_new calls the production entry gemm_prefill_smem_prewarm_one(tm, ks, af32) once per process for every launchable combination (the MINFER_GEMM_OPTIN_SET X-macro in src/cuda/kernels/common.cuh, shared with the fatbin lookup and the test seam). The placement is the argument: try_new runs under CUDA.get_or_init, before the state is published, before any CudaBackend exists, and therefore before the per-instance stream graph_begin_capture needs — so "the attribute is set outside any capture window" holds by construction, not by inference from the warmup count or the capture mode. It is once per process, not once per backend. The entry drives the same gemm_smem_optin<TM,KS,AF32> the launcher reads, so the pre-warm and the lazy path share one per-instantiation cache and one cudaFuncSetAttribute site; a cache-keying regression therefore cannot hide behind the pre-warm (it would leave the pre-warmed instantiations un-opted-in, which the gates read back from the device). A successful pre-warm prints nothing; a failure or a deliberate over-limit skip is named per instantiation, and MINFER_NO_GEMM_PREWARM=1 is the documented control that skips the loop (same-binary A/B and the "lazy path alone" gate arms).

The lazy per-launch opt-in stays as defence in depth: gemm_smem_optin (called from the launcher's GEMM_ONE) invokes the shared minfer_smem_optin helper on an instantiation's first launch and caches the answer in a function-local static per instantiation (a test injection is never cached). The helper names the site, the instantiation, the attribute, the requested bytes, the queried device limit and cudaGetErrorName, clears the latch, and the launcher does not launch when the answer is false; its own <<<>>> error is read too (minfer_launch_ok) and returned as 0, which prefill_gemm_f16_inner turns into an Err. The af32 wrapper launch_gemm_f32a returns the same result. A request above cudaDevAttrMaxSharedMemoryPerBlockOptin is skipped without calling the attribute on both paths, with the reason named. The same treatment covers every MMQ launcher (launch_mmq_nt/launch_mmq_raw_nt/launch_mmq_raw_nb_nt/launch_mmq_raw_nb_bt_nt/ launch_mmq_raw_nb_bt_q6k_nt/launch_mmq_raw_wide_nt); the terminal two return 0 → Err at the Rust caller, the fallback ones keep their documented 0 = clean fallback contract. The operator's signal stays the first-launch site report: a refused or skipped instantiation prints once, at the launch site (minfer_smem_optin). The removed checked/skipped counters did not come back, and a fully admitted pre-warm is silent.

Why the invariant holds — the pre-warm by construction, then three defence-in-depth mechanisms. The historical claim was that cudaFuncSetAttribute is illegal inside a capture window and poisons the context (error 700). Nothing in the repository establishes whether that holds for the adopted mode (see the honest limit below), so the design does not rely on the call being legal in a window; it relies on the call never happening in one. The eager pre-warm makes that true by construction (above). The three emergent mechanisms remain, now as the lazy path's fallback — called out at their sites (graph_replay_step in src/graph/cuda_backend.rs, gemm_smem_optin in src/cuda/kernels/gemm_wmma.cu); any change to one is a design change, not a tuning knob:

1. The 3-run capture warmup (capture_warmup, default 3): capture opens only from the third run of a (uid, range) key, and prefill-shaped graphs (nt > 1) capture by default since R3-B. An instantiation's first launch — the one that calls cudaFuncSetAttribute — therefore always runs uncaptured.
2. cudaStreamCaptureModeThreadLocal (#188's measured choice): the window belongs to the capturing thread, so a foreign thread's driver call cannot join it or invalidate it. Changing the mode must not silently change (1)'s guarantee.
3. The per-instantiation cache: the in-window launch re-reads the cached answer instead of asking the driver again, so the >48 KiB launch inside the window never calls the attribute. The cache is instantiated on the (tm, ks, af32) template parameters, not on a deduced K: every gemm_f16_nt_kernel_t shares one signature, and the pre-#218 template <typename K> gave the whole family one static (the #218 coverage gate found <64,64,true> answering for <128,64,false>, whose attribute had never been set). #223's pre-warm drives this same function, so the cache is exercised for every instantiation at context creation — the regression can no longer hide behind a sweep that bypassed the cache.

The #218 gates pin that as observed behaviour. cuda_prefill_smem_optin_is_done_by_production (a real prefill forward in a fresh process; asserts the device's own opted_in read-back), cuda_prefill_smem_optin_refusal_fails_the_prefill (the control arm: MINFER_TEST_CALL_FAIL=attr:gemm_f16_f16 makes the production prefill refuse the launch and name the site), cuda_prefill_smem_optin_is_never_set_inside_a_capture_window (a >48 KiB prefill captures, replays bitwise, and gemm_smem_optin_in_capture_count() == 0), and cuda_prefill_smem_lazy_optin_admits_every_launchable_instantiation (every launchable >48 KiB instantiation reads back opted in through the production function). The capture_warmup test seam (MINFER_TEST_CAPTURE_WARMUP=1) is the mutation lever for the counter.

#223 adds the runtime guarantee's own gate. issue223_tests::cuda_prefill_smem_prewarm_opts_in_every_launchable_instantiation_before_any_launch runs in a fresh process and asserts, immediately after CudaState::init() and before any kernel launch, that every launchable >48 KiB instantiation already reads back opted in; a second fresh process with MINFER_NO_GEMM_PREWARM=1 asserts the negation, so the read-back is not always 1. It is the detector for the mutation the #218 gates cannot see — a pre-warm that skips one (tm, ks, af32) (the lazy path simply opts it in on first launch, so the coverage/counter arms stay green). The four #218 arms run their fresh-process children with MINFER_NO_GEMM_PREWARM=1, the documented control, so their claims stay the lazy path's and the pre-warm cannot make them vacuous.

Cost. The pre-warm loop's own duration is the fatbin's one-time module load — which the process pays before the first kernel from it can run either way — so with MINFER_OP_TIMING=1 the line reads ~2.3 ms warm, and the pre-split/synthetic 152.9 µs proxy was not the cost. The full measurement is in docs/SOURCE-LAYOUT-PLAN.md §5.1: the pre-split and post-split per-module transcripts (warm and page-cache-cold — the correction that the ~6× cold factor is page-cache, not the GPU clock), the 17-module split's effect on it, and the MINFER_OP_TIMING reading that only looks smaller because it now times one seventeenth of the work. The Step 5 record in docs/ARCHITECTURE-EXECUTION-PLAN.md holds the ticket-level entry.

The net effect is still the proxy's conclusion: the module load moves rather than appears — but only because something later would pay it anyway. prewarm_prefill() (the r59 rider's minfer_prewarm_kernels, called at the end of Qwen2/Qwen3 weight registration, before the first forward) already pushes that fatbin load into the startup path; with the pre-warm on it costs ~2.3 ms, with the pre-warm off ~4.5 ms — the same ~2.2 ms, moved earlier. The coupling is explicit and load-bearing: the ≈ 0 net holds only while a later step pays that same module load — today prewarm_prefill(), otherwise the first launch from the fatbin. Move that rider after the first launch, or remove it, and the pre-warm's loop becomes ~2.2 ms of net-new startup cost of the same clock-dependent magnitude. Controlled probe on the real binary (fresh process, MINFER_MMQ=0/1 × MINFER_GEMM_K64=0/1): the first prefill forward is 2288–2314 µs with the pre-warm off and 78–120 µs with it on; prewarm_prefill() is 4.3–4.6 ms off vs 2.2–2.4 ms on. The hot path is untouched: minfer bench -p 2048 -n 128, same binary, 7 interleaved matched rounds, medians — pp2048 2546.55 vs 2544.34 t/s (+0.09%), tg128 236.30 vs 236.45 t/s (−0.06%), bar ±1%. The full transcript is in the #223 record.

The same-thread ThreadLocal in-window case — measured (2026-09-29). The open question this section used to carry — is cudaFuncSetAttribute legal inside a same-thread cudaStreamCaptureModeThreadLocal window? — is now measured: the MINFER_TEST_CAPTURE_WARMUP=1 mutation arm drives the opt-in into the first, captured run, and on GB10 sm_121 / CUDA 13.0 / driver 580.178.04 the call is tolerated — the attribute publishes, the

48 KiB graph still captures, instantiates and replays bitwise-identically, and only gemm_smem_optin_in_capture_count() moves. (The 2026-09-25 probe below measured the Global mode instead; this one is the adopted mode.) The design nevertheless keeps the call out of the window: that behaviour is not contractual across toolkits, and the counter gate is what makes "never set inside a window" an observed property rather than a driver assumption. The full transcript is in the #218 record.

Measured correction (2026-09-25, CUDA 13.0 / driver 580.178.04 / sm_121). A probe (/tmp/fix147_attr_capture_probe.cu) shows the historical claim above no longer holds verbatim on this runtime: cudaFuncSetAttribute returns cudaSuccess when called inside an open cudaStreamCaptureModeGlobal window, both for the already-set value and for a new one. The eager sweep that was in tree then has since been removed by #218 (it was already dead code — #188 had dropped its caller), and because the launcher caches one answer per instantiation it does not re-ask inside a window either.

#223 forward note (2026-09-29): the eager half is back, as a pre-warm through the lazy entry, not as the old sweep: CudaState::try_new calls gemm_prefill_smem_prewarm_one for every launchable instantiation, i.e. the same gemm_smem_optin cache the launcher reads. Nothing in the measured driver behaviour above changes; the point of the placement is to stop depending on it (the attribute is set before any window can exist, by construction). On the real binary the loop's measured cost is ~2.2 ms (the fatbin's one-time module load), which prewarm_prefill() already paid during registration — net new startup cost ≈ 0, hot path unchanged. See §2.4 and the #223 record in ARCHITECTURE-EXECUTION-PLAN.md.

2.5 Legacy surface

The pre-graph imperative path still compiles but is not driven by inference: layer_gpu, output_norm_gpu, init_kv_cache + the kv_k/kv_v/kv_size slots, the buf_* persistent-slot pool with get_or_grow/upload_*, the host-staged quant_matmul_* wrappers, and the old single-slot capture API. Each carries a scoped #[allow(dead_code)] with a reason; the release build is warning-free. The graph path owns KV through the allocator's persistent regions, so init_kv_cache is bypassed entirely.

3. llama.cpp Reference Map

llama.cpp's CUDA backend (ca3d5a3e1 at the time of the port) was the reference for the device layer, the replay state machine and the kernel strategy. What was borrowed and what was deliberately not:

The one llama.cpp idea still on the table as a step function is the q8_1 GEMM-prologue fusion (see CUDA_OPTIMIZATION.md §1.4); everything else in the campaign is closed or sub-bar.

4. Design

4.1 CudaBackend lifecycle and state

with_layout() resolves the device singleton, creates the instance's own non-blocking stream (issue #188; a device that cannot give one gives no backend), snapshots the engine's KV element type (crate::cuda::layout_of), reads the two graph gates (MINFER_NO_CUDA_GRAPH=1 → GraphMode::Disabled; prefill_capture defaults ON unless MINFER_NO_PREFILL_CAPTURE=1), and starts with an empty pool. Drop binds the stream, frees every pool pointer, the positions scratch, the cross-backend staging and every captured graph exec, then destroys the stream.

State groups:

Pool rules worth restating because they carry correctness weight:

• Exact byte-length reuse only — alloc_buffer scans free for pool[id].bytes == size * 4.
• free_buffer never frees — a persistent KV region must survive rebuilds, and the pool keeps device memory for the next graph. Only Drop returns memory to the driver.
• alloc_fresh exists for split-boundary staging: ids in the free list are still referenced by node_to_buf and physically live during the execute that follows.
• OOM is not a panic. cuda_malloc logs and returns null; the null buffer fails cleanly at execute time (ptr_of). Panicking is forbidden because it would poison the shared scratch maps and the device-entry token (the legacy path) for every other user.

4.2 Backend trait mapping

4.3 Eligibility

supports_op is f32-activation only (dtype != DType::F32 ⇒ false) and answers:

• Unconditionally supported: Input, Add, Mul, Silu, SwiGLU, RmsNorm, QkNorm, MatMul, Attn, KvcacheStore, KvcacheLoad, View, Reshape, Permute, GetRows, FusedQKV, QkvBiasRopeStore, FusedFFN.
• Conditional: RoPE only for RopeStyle::NonInterleaved (the neox layout; the only style the supported architectures emit).
• Everything else (Scale, Softmax, BatchMatMul, FusedQkvNorm) stays on CPU.

Two layers of checks are deliberately not in supports_op:

1. Weight/quant eligibility is a model-level, all-or-nothing gate (weights_on_cuda, §4.6). A layer whose weights are not all registered in a kernel-supported type keeps the whole graph on CPU rather than creating a partial-GPU split. The whitelist for matmul weights is Q4_0/Q4_1/Q5_0/Q5_1/Q8_0/Q4_K/Q5_K/Q6_K/F32; the embedding (tok_embd) additionally has its own type list because it is gathered, not multiplied.
2. Shape and feature invariants are enforced in execute_node and return Err — never a silent fallback. Guards include: transpose_b unsupported; quantized matmul id % 32 == 0; RmsNorm/ QkNorm dim a nonzero multiple of 4 (the float4 kernel); attention nkt == n_head_kv * hd, hd == hd_kv, hd a multiple of 4 in 1..=128, n_head % n_head_kv == 0; the fused decode nodes nt == 1; RoPE neox; KvcacheStore's output buffer must be the K region; a missing declared norm weight is an error.

4.4 Execution dispatch

execute_node_inner opens with one stream-lock acquisition and a cache rule:

The MMQ A-quantize memo is valid only across consecutive MatMul/FusedFFN nodes; every other node kind clears it (clear_mmq_cache). FusedFFN is in the preserve set because its input is the FFN-norm output whose pre-quantized plane the fused producer just recorded.

MatMul family selection (matmul_f32_ptr_layout in cuda.rs):

• Prefill GEMM when (nt >= 9 || MINFER_SMALL_M_GEMM=1) && id % 32 == 0 && !MINFER_NO_PREFILL_GEMM and the type is quantized: the promoted int8 MMQ path (MINFER_MMQ, default on) — pad40 pre-transposed A planes, fused producers, raw-byte NB-BT tensor-core kernels for q4_K/q6_K; else the f16 wmma GEMM (or the persistent f16 weight cache / MINFER_FUSED_B dequant-in-GEMM).
• Decode / small batch: per-type MMVQ (dp4a over q8_0 activations) for nt == 1 with shape gates, _multi variants for nt 2..8, and f32-activation kernels otherwise. Q4_1/Q5_0/Q5_1 have f32 kernels only; F32 weights use f32_f32_matmul_vec/_scalar.
• f16 weights (#141): f16_f32_matmul_vec when id % 8 == 0, else f16_f32_matmul_scalar — for every nt, because f16 is not an MMQ format (MMQ streams quantized bytes) and the f16-wmma path is the MINFER_MMQ=0 fallback for the quantized types, not an f16-weight kernel. The weight bytes stay 2 B/element on the device: there is no registration-time dequant to f32, so the memory the f16 file exists to save is actually saved (a 0.5B f16 GGUF registers 942.4 MiB of device weights; ~1.9 GiB if it were dequantized). The kernel converts in-register with __half22float2 and FMA's against the f32 activations, so the accumulation is f32 like every other CUDA matmul. Both launchers read their own launch return through the #147 helpers and return non-zero → Err, rather than joining the unchecked <<<>>> sites of #162. The f16 embedding gather is embed_rows_f16 (one thread per output element) — without it the all-or-nothing weights_on_cuda check would drop a converted f16 GGUF to the CPU over its token_embd alone.
• bf16 weights (#208, the CUDA half): the exact sibling of the f16 pair — bf16_f32_matmul_vec when id % 8 == 0, else bf16_f32_matmul_scalar, and embed_rows_bf16 for the embedding. The promotion is cheaper than f16's: bf16 is f32's top 16 bits, so the vec kernel loads 8 elements as one uint4 and shifts each word (b2f(bits) == __uint_as_float(bits << 16), exact — no rounding at all, unlike the quantized types), while the scalar kernel shifts one word per element. The launchers are their own (launch_bf16_f32_matmul / launch_embed_rows_bf16), each reading its own launch return through the #147 helpers, so a failed launch is an Err at the call site. The same two design decisions as f16 apply and for the same reasons: the weights stay 2 B/element on the device (the f16 bullet above carries the measured figure — the f16 twin's number — where a dequantized copy would be ~1.9 GiB), and a bf16 prefill does not enter the int8 MMQ GEMM quantized types only; MMQ streams quantized bytes and bf16 is not a format). bf16 is registered by the shared models::weight_reg::cuda_weight_reg rule, so both architectures admit it in one place; cuda::concat_rows has no 2 B/element arm, so the attn_qkv / ffn_gu concat copies are not registered and the bf16 graph runs the unfused matmul chain. The exactness gate is cuda_backend::tests::weights::cuda_bf16_matmul_matches_the_exact_shift_reference (both launcher arms, bitwise against f32::from_bits(bits << 16)) and cuda_backend::tests::weights::cuda_bf16_embed_gather_matches_the_exact_shift_reference; the placement/real-model gate is f208_bf16_weights_run_on_the_cuda_device (169 bf16 matmul + 1 embed nodes all on CUDA; device vs CPU max |Δlogit| 7.82e-5 / 4.24e-6 relative, bar 0.01 / 1e-3, greedy identical). The Metal twin is #208's other half and landed separately (PR #323), which widened the Metal registration arm to matches!(ttype, F32 | F16 | BF16) — see METAL-BACKEND-DESIGN.md §4.4.
• Bias is applied by add_bias_f32 after the GEMM; its last argument is the row count nt (the kernel maps one block row per token).

4.5 Allocator and scheduler integration

• GraphAllocator::supports priority is Metal → CUDA → CPU; enable_cuda() mirrors enable_metal(). On a CUDA-only host the practical effect is CUDA first.
• KV compaction (C3) is not a node op — it is the Backend::copy_cells trait method, called by GraphAllocator::kv_defrag between forwards. CUDA implements it with kv_move_rows: one block, rows walked ascending when dst_row <= src_row and descending otherwise (C7b — a compaction slides a run down into the gap below it, and growing a run can move one up), with a __syncthreads() between rows, because the ranges may overlap in either direction and device-to-device cudaMemcpyAsync is documented undefined for overlap. No staging buffer, no second pass. The launcher returns non-zero on a contract violation and the Rust side turns that into an Err, so the allocator fails the compaction before it renumbers any run. It runs on the backend's own stream, so it is ordered after the previous forward's kernels.
• KV regions are created by ensure_kv on the layer's assigned backend, so with CUDA assignment the per-layer K/V regions live in the CUDA pool and KvProvider::kv_pair returns pool ids that execute_node resolves to device pointers. init_kv_cache is bypassed.
• Cross-backend values go through the allocator's staging path: copy_to_cpu (pinned D2H) + write_host (pinned H2D) into a buffer allocated with alloc_fresh on the consumer's backend.
• The scheduler syncs at split boundaries (sync_backend → CudaState::sync()), which is also where a pending capture window closes. In practice the post-7e③ graph is a single CUDA split for Qwen2/Qwen3 (embed gather and tail gather are on device), so there are no per-step cross-backend copies at all; splits appear only in synthetic or intentionally mixed graphs.

4.6 Model wiring

Both architectures use the same shape (Qwen2 shown; Qwen3 mirrors it):

cuda_on  = CudaState::get().is_some() && weights_on_cuda(model)     // #[cfg(feature = "cuda")]
metal_on = metal_available() && weights_on_gpu(model)
CParams.gpu = metal_on || cuda_on

• weights_on_cuda is the all-or-nothing gate: every graph-referenced tensor must be registered (has_weight_of_size) and, for matmul/embedding tensors, of a kernel-supported type. On failure it prints CUDA GATE: weight '<name>' (type <t>) has no CUDA kernel or is not registered and the model runs entirely on CPU. Since 7e③ tok_embd is gated like every other weight (the embedding gather is on device); its own type list is F32/Q4_0/Q8_0/Q4_K/Q5_0/Q5_1/Q6_K/Q5_K.
• Registration happens in the loader, not in register_graph_weights (which fills the CPU registry only): cuda.register_weight per tensor, register_weight_q6k_padded for Q6_K, register_weight_q80_p32 for the q8_0 split plane, and the _exp/_dsc/__dpl planes under their gates. The loaders also build the fused concat weights with cuda::concat_rows and register them: blk.{i}.attn_qkv (Qwen2 only — the Qwen3 fused-QKV path is Metal-only, so CUDA registers no Qwen3 attn_qkv) and blk.{i}.ffn_gu (both models, gated on nf <= 16384).
• Concat availability: Qwen2's qkv_concat_available / gu_concat_available use the metadata-only cuda::concat_rows_feasible probe on the CUDA arm, because rebuilding the concat bytes during graph construction measured a ~920 ms stall per decode graph build. Qwen3's gu_concat_available uses the eager cuda::concat_rows (its concat is built once per layer at load either way); Qwen3's qkv_concat_available has no CUDA arm.
• Fused nodes per model: Qwen2 builds FusedQKV (concat class) or QkvBiasRopeStore (mixed-quant class, CUDA-only) for decode QKV, and FusedFFN when nf <= 16384. Qwen3 builds FusedFFN but not its FusedQkvNorm on CUDA: that fused path is Metal-only (the Qwen3 concat probe has no CUDA arm), so CUDA Qwen3 runs the unfused qk_norm chain.
• FusionPass receives the enabled backends in its Vec<&dyn Backend> (F4: GraphAllocator::fusion_backends) and its node → index map is GraphAllocator::fusion_backend_index, so the SwiGLU rewrite is gated by CUDA's own supports_fused and the two cannot drift apart. The pre-F4 hand-built vector and its b.name() == "cuda" position lookup are gone.
• The dump/debug tags in forward_cached are backend-agnostic (MINFER_GRAPH_DUMP; MINFER_REBUILD_TRACE=1, Qwen2 only). MINFER_CUDA_DEBUG was a device-layer trace on the legacy layer_gpu surface; #240/#241 deleted that surface, so the knob and its per-node syncs (debug_sync) are gone — the graph path's CudaState::sync() is the remaining drain point.

4.7 Memory, residency and staging

Residency. Weights are uploaded at load and never copied during execution; only activations, positions and logits cross PCIe (tiny for decode). Quant-specific registrations trade device memory for speed: the Q6_K padded repack (224-byte slots), the q8_0 p32 split plane (≈ +94% of that tensor's bytes), the q6_K W_exp dense plane (~1.52 GB on 7B), the q4_K/q6_K W_dsc f32-pair planes (~1.46 GB), and the D4-4 q6_K dense split plane. Each has an opt-out gate (§5, env-gate reference in CUDA_OPTIMIZATION.md Appendix A); the promoted default spends ~3.27 GB to reach the 1.080× prefill path.

KV layout. The persistent K/V regions keep their f32 IR shape, but the store/attention kernels run one of three layouts, tagged by crate::cuda::KV_LAYOUT_F32/F16/Q8_0 — the same 0/1/2 codes KvFormat uses, and a host contract the kernels are templated on (int LAYOUT in the src/cuda/kernels/*.cu KV/attention TUs):

• KV_LAYOUT_F32 — one f32 per element;
• KV_LAYOUT_F16 — one f16 per element in the first half of the f32-shaped region (kvformat::auto_device_format selects it when n_layers × n_kv_embd >= 8192, the 7B class, and MINFER_CACHE_TYPE=f16 overrides). Its snapshots are restorable since #130: the session container records the type in its header flags (FLAG_F16, C5 S3), so an auto-f16 model's --slots-file / --session companion is no longer refused on load;
• KV_LAYOUT_Q8_0 — packed 34-byte Q8_0 blocks, one cell rounded up to whole f32 words (MINFER_CACHE_TYPE=q8_0, C4 S2b).

Every KV address is formed in bytes: kv_row(base, cell, row_bytes) names a cell and kv4<LAYOUT>(row, elem) -> float4 is the one load idiom — the old float4 load for f32, the old two-__half2 pair for f16 (both bit-identical to the pre-C4 instantiations), and for Q8_0 the f16 scale plus four quants of block elem/32. A 4-element group never straddles a block because a KV head's base is hd-aligned and hd % 32 == 0 (ensure_kv's packed-width check). The kernels take const void* k/v plus size_t row_bytes, and the launchers take the layout as an int.

CudaBackend holds that int in its kv_layout field, and kv_row_bytes(nkt) derives the stride (nkt*4, nkt*2, or KvFormat::Q8_0.row_bytes(nkt)). The packed store is store_kv_q8_0, whose quantizer is the CPU's step for step (amax/127, f16 scale, round-ties-even), so both backends store the same bytes. Before C4 S2b this was a bool that mapped anything not exactly f16 to f32 — which would have addressed a packed region as f32 rows, the silent corruption the layout tag exists to make impossible.

Per-engine scope (#99, completed by #153). #99 made the KV format per engine for the model, the graph builder (CParams::kv_format), the allocator and the CPU kernels; #153 finished the device half. There is no cuda::KV_LAYOUT static any more: the engine's resolved KvFormat — including the GPU auto policy, which kvformat::resolve now folds in from the model dims — reaches CudaBackend through GraphAllocator::set_kv_format (and enable_cuda builds a fresh backend from the same stamp), so CudaBackend::kv_layout is the only source the dispatch reads. cuda::layout_of / format_of are the one binding between KvFormat and the FFI tag. The launchers already took the tag as an argument; what was process-wide was the value. Two engines in one process therefore run their own layouts, and the captured-graph key carries the tag (below), so an exec instantiated for one layout cannot replay for another. models::load_model_configured no longer restates anything.

All three backends keep the format per engine. Metal joined in #44 part (b); the decision and the deleted process-wide symbols are in ADR-0005 and ADR-0006, and the Metal detail is docs/METAL-BACKEND-DESIGN.md.

Host transfers.

• H2D fills go through a lazy ring of 8 × 2 MiB pinned slots (write_input_async): the Rust slice is copied into a slot and the cudaMemcpyAsync is queued; same-stream ordering guarantees the consumer kernels see the data. A ring wrap retires in-flight copies with one stream sync; oversized inputs or a failed cudaHostAlloc fall back to a blocking copy.
• D2H readbacks use a single grow-on-demand pinned buffer (PinnedBuf), pre-grown to 4 MiB at first prefill because the lazy cudaHostAlloc showed up as a 0.78 ms tail malloc at the logits readback. MINFER_NO_PINNED_READBACK=1 reverts to pageable copies.
• Device memory is not host-readable by plain memcpy on GB10 — a host probe that dereferences a device pointer faults. All D2H goes through cudaMemcpy staging (copy_to_host).
• Trace/viz uses CaptureStaging: one pinned buffer (128 MiB ceiling) into which each captured node's output is queued as a stream-ordered async D2H right after its launch, drained with a single sync at the split boundary. Node outputs above the ceiling fall back to the per-node synchronous copy.

Caches. The MMQ A-quantize memo (consecutive-window rule above), the positions→i32 memo (keyed by (input buffer id, pool_gen), cleared in synchronize), and the optional persistent f16 weight cache (w16_cache, enabled when quantized matmul weights ≥ 2 GiB and MMQ is off) are all execution-window caches with explicit invalidation points rather than long-lived state.

4.8 CUDA Graph capture and replay

graph_replay(uid, range, nt_hint) is called by the scheduler once per split, before its node loop. The state machine:

1. graphs_mode != Enabled → direct launches.
2. An open capture window of our own → direct launches (a nested replay would be CUDA-invalid; the graph is single-split today so this is unreachable).
3. A stored exec for (uid, range) with a matching pool_gen and kv_layout → cudaGraphLaunch; on launch failure, disable graphs for the session and fall back.
4. A stored exec with a different pool_gen (pointer layout may differ) or a different kv_layout (the recorded kernels were instantiated for the old tag — store_kv_f16 vs store_kv_q8_0, the layout-tagged attention) → destroy it, drop the warmup counter, re-warm. #153 added the kv_layout term; CudaBackend::set_kv_layout also invalidates eagerly when a stamp moves, so the lookup check is the second line of defence.
5. Warmup: executions 1 and 2 of a key run direct launches (llama.cpp warms up twice); a one-shot prefill never reaches capture.
6. On the third execution, if nt_hint.map_or(true, |nt| nt == 1 || prefill_capture), the backend opens a capture window on its own stream, in cudaStreamCaptureModeThreadLocal (§2.4) — no process-wide lock, because no other backend shares this stream. The gate means decode-shaped graphs always capture; prefill-shaped graphs capture only when prefill_capture is on (default ON since R3-B; MINFER_NO_PREFILL_CAPTURE=1 opts out).
7. The window closes at the split's synchronize → close_capture_or_sync: end capture, instantiate, launch once so the step still produces output, cache the exec at the current pool_gen and kv_layout, then sync. A failure destroys the exec, logs loudly, and disables graphs for the session.

Replay correctness rests on stable addresses: pool ids never move memory, copy_across rewrites the same staging buffers each step, and the positions scratch pointer is embedded in captured execs — so growing it bumps pool_gen and forces re-capture.

Two interactions are part of the contract:

• Trace/viz disables replay. The scheduler skips graph_replay entirely while MINFER_TRACE or live viz capture is active, because per-node host readbacks inside a capture window are illegal.
• A node error inside the window aborts it (abort_capture): end capture without launching, destroy the exec, disable graphs, sync. Later steps run direct-launch with graphs disabled; the aborted step's outputs were never produced and are consumed as-is — there is no poisoned-error mechanism, and the code says so explicitly.

4.9 GPU safety (CUDA edition)

The hard rules live in docs/GPU_SAFETY.md (CUDA section); this is how the backend implements them:

1. Errors are errors. Guards return Err naming the node and the blocking values; the scheduler aborts. There is no mid-run CPU fallback.
2. No sync inside an active capture window. The 7e② incident — a temporary sync wrapper that produced garbage only with graphs on — is the recorded reason. Debug reads go through the boundary.
3. D2H always through staging (GB10 device memory is not host-readable by plain memcpy).
4. A latched error is never blamed on the kernel that just ran. sync() polls cudaGetLastError + cudaStreamSynchronize; the first reports whatever an earlier call on the thread latched, so its message names the observer and the cudaGetErrorName symbol and says it is not attributed to a kernel (the error is counted and cleared, not dropped). The two origins behind the old phantom "kernel launch error: 1" were a rejected cudaFuncSetAttribute and cudaGraphDestroy called on a cudaGraphExec_t — both fixed at their call sites by #145, and the remaining unchecked sites of the same class (every MMQ dynamic-smem opt-in and launch, the prefill-GEMM launcher's own launch, and graph_end_capture_to_exec's cudaGraphDestroy) by #147. #162 then removed the class entirely: every <<<>>> in src/cuda/kernels/*.cu reads its own error, so a latched error at sync() is by construction an error no site read (a non-launch API call), never an unattributed launch.
5. A return value that gates a later launch is read where the call is made. Every dynamic-smem opt-in goes through minfer_smem_optin: an over-limit request is skipped with the reason, any other failure is named and cleared at the call site and the launch is refused (a launch over an un-opted-in dynamic smem cannot succeed). Every launch goes through minfer_launch_ok, whose immediately-following cudaGetLastError is a launch check because minfer_launch_prelude cleared (and reported) any latch that predates the launch. The sync poll is the backstop for a missed site, not the place to diagnose one; #147's issue147_tests gates inject a real failure at every one of the sites and assert the named report, the refusal and a clean latch. #162 states the per-op severity in the helper (the plan's #162 record holds the site census): minfer_launch_ok is required — it records a sticky failure that CudaBackend::execute_node turns into an Err naming the site (one Rust-side check, not one per launcher, so the op never proceeds on a stale output) — while minfer_launch_ok_opt only names and clears for a path with a documented fallback (the MMQ fast paths, the fa-prefill smem fallback, and the int-returning launchers whose Rust caller already decides). Both levers are data: minfer_launch_block (an illegal block geometry) and minfer_launch_smem (an over-limit dynamic smem request) make the real launch fail.
6. Same-stream ordering is the async-fill contract; the pinned ring syncs on wrap and never hands a slot back early.
7. Weight-registry ownership: name+size reuse, different-size replace with a deliberate, bounded leak (a live captured graph may still reference the old buffer).
8. Device limits are queried at runtime (SM count, compute capability, free memory); the only hardcoded shape knowledge is the compiled target list and the documented kernel invariants.

5. Implementation Phases

5.1 Phase 7 (7a–7e)

5.2 After Phase 7

The optimization campaign is indexed in docs/CUDA_OPTIMIZATION.md §0 and expanded in docs/cuda_optimization_steps/. For the design record, its eras and what each one changed in the backend or its graph integration:

The retired CUDA-FOLLOWUP-PLAN.md was consolidated into step docs 78/79 (commit ace6242); its residual open items are listed in CUDA_OPTIMIZATION.md §1.4.

6. Risks and Open Questions

The plan's original risk table, with its resolution:

Newer, still-standing items:

7. Verification

7.1 Device test suite

src/graph/cuda_backend.rs carries the device suite (39 test functions); the device layer's own gates live in src/cuda.rs. Both are device-gated: tests skip when no CUDA device is present. Categories and the invariants they pin:

• Capture / replay — cuda_graph_replay_bit_parity, cuda_prefill_shaped_graph_never_captures, cuda_multisplit_capture_bit_parity, cuda_prefill_capture_defaults_on, cuda_prefill_capture_bit_parity_pp16_pp300, cuda_capture_abort_on_error, cuda_graph_recaptures_on_pool_gen_change, cuda_graph_generation_replay_parity_real_model, cuda_capture_staging_order_and_fallback; the device-layer gate cuda_graph_exec_destroy_leaves_no_latched_error (#145: destroying the cudaGraphExec_t must not latch an API error) and the env-gated cuda_sync_surfaces_a_latched_error_as_latched.
• Pool / allocator / host transfer — cuda_pool_roundtrip, cuda_pinned_readback_roundtrip, copy_across_cpu_to_cuda_and_back, kv_persistent_regions_survive_realloc, cuda_scheduler_chain, cuda_row_marginal_bench (env-gated bench).
• Per-op parity — elementwise (cuda_elementwise_parity), norms (cuda_norm_parity), matmuls (cuda_matmul_parity, cuda_kquant_matmul_parity, cuda_q5_matmul_parity, and the per-type MMVQ parity tests for q4_K/q6_K/q5_K), attention (cuda_attn_split_decode_parity, cuda_verify_attention_nt_invariance, cuda_rope_kv_attn_roundtrip, cuda_kv_f16_roundtrip_attn), embedding gather (cuda_embed_getrows_parity), fused FFN (cuda_fused_ffn_parity).
• KV cell movement (C3/C7b) — cuda_copy_cells_moves_overlapping_rows_in_both_directions (rows [1, 4) -> [0, 3) and then [0, 3) -> [1, 4), i.e. two of three rows are read and overwritten in each direction; the whole buffer is compared against the bytes a copy through a temporary would produce. The kernel walks the rows in the order the overlap requires — ascending when the run slides down, descending when it slides up — and kvcache::order_moves fixes that order across several runs).
• Prefill GEMM / MMQ / FA — cuda_prefill_mmq_parity, cuda_prefill_f16_gemm_parity, cuda_q4_0_prefill_q8_0_gemm_parity, cuda_fa_prefill_attention_parity (note: causal, single-sequence, start = 0 — the windowed FA mask is covered by cuda_windowed_attention_matches_causal_for_long_windows, which is what caught the fa_prefill_f16kv window-limit fault on 2026-09-19), the byte-exact plane tests (cuda_q6k_exp_dense_byte_exact, cuda_q6k_dsc_dense_byte_exact, cuda_q4k_dsc_dense_byte_exact), cuda_prefill_fused_b_bitparity, and cuda_multi_token_matmul_bitwise (one nt=3 forward bitwise-equal to three nt=1 forwards). The lazy >48 KiB opt-in (#218) has its own gates: the_gemm_smem_formula_matches_the_kernel_layout (pins gemm_dynamic_smem_bytes against the kernel's byte layout — the value-level arm, so a silent shrink is seen); cuda_prefill_smem_optin_is_done_by_production (a real prefill forward makes the device read back opted_in == 1 for gemm_f16_nt_kernel_t<128,64,false>, in a fresh process so the "not opted in before" precondition is observable); the control arm cuda_prefill_smem_optin_refusal_fails_the_prefill (MINFER_TEST_CALL_FAIL=attr:gemm_f16_f16 refuses the launch, and the site report names the call and cudaErrorInvalidValue); cuda_prefill_smem_lazy_optin_admits_every_launchable_instantiation (every launchable >48 KiB instantiation reads back opted in through production's own gemm_smem_optin, and every over-limit one is refused without a call); and cuda_prefill_smem_optin_is_never_set_inside_a_capture_window (a >48 KiB prefill captures, replays bitwise, and gemm_smem_optin_in_capture_count() == 0 proves the opt-in ran before the window opened). Plus the_latched_error_message_never_blames_a_kernel. #147's issue147_tests module adds the_graph_destroy_failure_message_names_the_matching_destructor (pure; the injection matcher itself moved to testfail::tests::the_matcher_is_exact_and_comma_separated in #171) plus the three cuda_issue147_* deliberate-failure gates (device; env-gated behind MINFER_TEST_ISSUE147=1, which arms MINFER_TEST_CALL_FAIL per site): every dynamic-smem site names its failed opt-in and refuses the launch, every launch site names its failed <<<>>> and refuses, and the destroy site names a failed cudaGraphDestroy and leaves no latch.

Model-level CUDA coverage: cuda_conversation_multiturn_reuse (Qwen2) asserts that an incremental multi-turn session reusing the decode graph and appended KV produces the same turn-2 text as a fresh conversation that re-prefills. graph_logits_match_forward_real_model's comparison helper switches to greedy-token equality when a CUDA device is present (the 7e① path-identity artifact). The Metal-only tests (fused_qkv_matches_unfused_decode, fused_qkv_norm_matches_unfused_decode, graph_metal_*) skip on a Linux CUDA build.

7.2 Acceptance gates

The campaign's verification methodology — the five-gate chain (parity ×3, greedy-32 identity, interleaved A/B medians with a +1.5% whole-prefill bar, the device suite, the ncu/nsys/SASS protocol) and its transferable lessons — is docs/cuda_optimization_steps/77-verification-methodology.md. Read it before running any A/B on this engine. Two standing measurement rules: never quote llama-bench long-context throughput as an attention target without an ncu byte-count or a llama-cli recall cross-check, and an end-to-end max|Δlogits| ≈ 0.38/0.39 is the inherent class of any accumulation-order change — argmax + greedy divergence + A/B are the operative gates.

Suite baselines in the campaign records: the Phase-7e entries say "144/0 (CUDA parallel + single), 130/0 plain"; the decode campaign's later rows reach 187/0/3. The non-CUDA baseline on the current tree is 145 passed / 0 failed / 3 ignored. Run cargo test --release for the plain suite and cargo test --release --features cuda on a device.

7.2a Profiling on dgxspark (ncu / nsys)

The reusable recipe, so the next session does not re-derive it (recorded 2026-09-27, GB10 sm_121, CUDA 13.0, driver 580.178.04):

• ncu is installed but not on PATH. The binary is /usr/local/cuda-13.0/bin/ncu (2025.3.1). which ncu finding nothing means the directory is not on PATH, not that the tool is missing — use the absolute path.
• A normal user cannot collect counters here — RmProfilingAdminOnly: 1 makes ncu fail with ERR_NVGPUCTRPERM. sudo collects (passwordless on dgxspark); no module parameter change and no driver reload is needed. Record "collected as root; module parameter unchanged".
• sudo changes HOME to /root, so the model cache under /home/yusiwen/.cache/minfer/models is invisible to a sudo-launched binary. Pass absolute model paths (MINFER_BATCH_TEST_MODEL=/home/yusiwen/.cache/..., or the path on the command line) or use sudo -E. A "model not found" under sudo is this, not a missing file.
• Build as the normal user first, then attach ncu to an already-built binary. ncu does not write repository files, so target/ stays yusiwen-owned; if a sudo run does leave a root-owned file, chown it back before the next cargo build.
• Some metrics are n/a on GB10/sm_121 — dram__bytes.sum among them (the integrated-memory / DGX Spark form exposes no classic DRAM counters). Confirm a metric name exists with ncu --query-metrics first, prefer the SM / instruction / L1 / L2 families, and treat a metric you actually collected as the only evidence; n/a is not a number.
• Collect targeted, not whole-run. Counter replay is slow: filter with -k regex:<kernel> and bound it with --launch-count, rather than replaying a full bench.
• A -k regex must match ncu's base kernel name, not the demangled signature. ncu lists the available kernels by base name (gqa_attn_split_partial, no template arguments), so a regex that includes < — e.g. -k 'regex:gqa_attn_split_partial<' — matches nothing and ncu prints an "Available Kernels" list instead of collecting (No kernels were profiled). Anchor it (-k 'regex:gqa_attn_split_partial$') so sibling kernels (..._combine, ..._bt) are excluded rather than eating into --launch-count. Recorded by #202, which lost one collection to this.

nsys stays the cheap, always-available instrument for per-kernel durations (nsys profile --trace=cuda --cuda-graph-trace=node ... — without node, kernels launched from a replayed CUDA graph are traced as one graph and never appear individually, which silently hides the whole decode path).

7.3 Recorded verifications — #145, #147, #141, #165, #167

These five are per-ticket verification records, and their durable home is the plan's ticket section (docs/ARCHITECTURE-EXECUTION-PLAN.md, one section per ticket). What their gates assert is stated where it belongs, not here: the failure-injection seam in docs/GATE-CONTRACT.md (#171), the launch-return rules in §4.9, and the q4_K dsc admission in §2.3.

The old §7.4–§7.7 and §7.9 headings were folded into this one; §7.8 and §7.10 keep their numbers because their records are not wholly duplicated — they carry the norm-weight invariant and the #189 value arm respectively.

7.8 The norm weight type is part of the rms_norm invariant

norm_weight(node, elems) resolves CudaState::weight_size(name) (the registered raw length, the same convention has_weight_of_size uses for a padded Q6_K plane) and requires exactly elems * 4, returning Err that names the registered length, the length the kernel reads, the node and "f16-norm". Both callers pass the dim the kernel uses (node.out_shape[0] for Op::RmsNorm, *hd for Op::QkNorm), so the check is the kernel's own geometry, not a name heuristic. It exists because the CUDA rms_norm kernel indexes the weight as d f32 elements regardless, so an f16 (2 B/element) norm weight — which minfer quantize --type f16 used to write for every 1-D tensor — would be read past its end.

The gate is device-only (CI has no GPU) and its arithmetic has no CI-covered pure twin: it is one usize comparison against elems * 4. It does not make an f16-norm model loadable on CUDA — registration still admits an f16 1-D weight, and the refusal happens at execute time with the node named (the Err-not-fallback rule of docs/GPU_SAFETY.md). The verification record (the mutation that made the f16 arm execute, the compute-sanitizer output, the suite counts and the #169 end-to-end acceptance run) is in the plan's #169 section.

7.10 The S4 map-window A/B: a value arm and a paired sign test

The defect. The gate was a pure stopwatch: it interleaved matched rounds of the span window and the map window and asserted the median of 9 per-round ratios <= 1.25x, which flips once five pairs are disturbed — and one parallel #[ignore]d device run had exactly five above the bar. The verdict was about how loaded the box was, not about the kernel. The failing transcript, the co-tenant reading and the bar's #123 provenance are in the plan's #189 section.

The fix, two halves.

1. A value arm first (gate contract rule 1). Before any timing the gate asserts
   • a one-row map window at cell 512 returns exactly that row's V — an absolute value computed on the host, not a relation between the two modes (softmax over one key is exactly 1.0, so the kernel is the identity on V);
   • a two-run map window at cell 512 returns the span's bytes over the same rows, bit for bit — f32 KV at the decode shape (nt = 1), f16 KV at the FA-prefill shape (nt = 512). One run is indistinguishable to a resolver that reads (cell, len) as (lo, hi); two runs, at a non-zero base, are not;
   • the map instantiation actually ran, by counted observation rather than the dispatch's own report: testfail::note_checked("cuda_attn_map_window") is bumped in the launcher (CudaState::gqa_attn_split / gqa_attn_kv_prefill, src/cuda.rs), and the gate resets it, runs one span call (counter stays 0) and one map call (counter 1). The bitwise arms say what was computed; the counter says the map path was the thing that computed it.
2. A paired sign test (gate contract rule 4). The timing verdict is the count of matched pairs whose map arm is above 1.25x its own span arm, and the gate refuses only at 7 of 9 — the one-sided sign test at alpha = 46/512 = 0.090. A load spike that disturbs a minority (or even a bare majority) of pairs cannot decide it; a doubled map cost moves all 9 and does. The bar is unchanged at 1.25x and the timed fixture is the pre-#189 one, so the recorded margins stay comparable; the round count is fixed in advance (PAIRS = 9) and every per-round ratio and the refusal count are printed. The reproducible mutation is MINFER_S4_AB_MAP_REPS=2 (src/cuda.rs::s4_ab_map_reps), which issues every map-mode attention launch twice — an implementation mutation, not a test edit.

Honest scope. Every number is a local GB10 measurement; CI has no GPU, so its CUDA job only compiles the harness. The two-run value arms are the gate's own fixtures, not the graph path — the graph-path bitwise coverage stays with cuda_map_window_matches_the_span_over_the_same_rows (which sweeps f32/f16/q8_0 over one/two/three runs and both batch shapes). The MINFER_S4_AB_MAP_REPS seam is wired to the two launches this gate drives (gqa_attn_split, gqa_attn_kv_prefill), not to the batched split path. The bar 1.25x is inherited unchanged from #123; this ticket changed the statistic and added the value arm, and did not widen it. The measured tables — the idle gate and six parallel runs, with every per-round ratio — the mutation output, the suite counts and the superseded src/device_entry.rs guard note are in the plan's #189 section.

8. Out of Scope / Future

• Not planned (revisit with a concrete need): cuBLAS/cublasLt, VMM pool, multi-GPU + peer copies, graph_optimize-style node reordering, Windows, self-hosted CUDA CI, IQ/Q2/Q3 quants.
• Open leads, all sub-bar or step-function (details in CUDA_OPTIMIZATION.md §1.4): the q8_1 GEMM-prologue fusion (the identified step change); the rms_nw roofline (+0.5–1%); a wave re-tile for small-od classes (+0.3–0.8%, needs ≤85 registers); a fused ffn_gu concat (needs the G5 nf <= 16384 gate re-measured); FA deep-opt only with a numerics-order-preserving structure.
• Inherited Phase-8 ledger: the shape-dependent q6_K MMVQ idle-tail rule (not started), the CUDA CI runner (deferred), and the /tmp/minfer_phase7/ ledger cleanup (awaiting a decision).
• Landed since the original plan and therefore no longer in this list: prefill CUDA-Graph capture, FA-style prefill attention, fused decode nodes on CUDA, pinned host buffers, f16 KV, Q5_K/Q5_1/Q5_0 and F32 kernels.

Running on a real device — stream ownership and launch checks (moved from AGENTS.md)

scripts/cuda_test.sh (cargo test --release --features cuda -- --test-threads=1) is the only way to run the device-gated tests — CI has no GPU and its CUDA job only compiles the harness — and because the device state is a process-wide singleton (CudaState) it must run serially (#64); a local run must rebuild the CLI with the feature (a plain cargo test --release overwrites target/release/minfer with a CPU-only build and silently measures the CPU), and MINFER_DISABLE_CUDA is presence-checked (=0 disables CUDA). The per-instance stream, capture-window and scratch design, the #188 probe and its measured table, the launch-return audit (scripts/check_cuda_launch_returns.py) and the injection lever (MINFER_TEST_ISSUE162=1, which drives every audited launch return through a real failing launch) are §2.4; the retired src/device_entry.rs guard is the #240/#241 record in docs/ARCHITECTURE-EXECUTION-PLAN.md.

Decisions governing this document

This page is the current contract; the decisions behind it are frozen in the ADR corpus:

• ADR-0001 — Inference runs through one declarative compute graph
• ADR-0002 — Topology is a function of GraphParams alone, so positions cannot be structure
• ADR-0006 — The KV storage format is a per-engine gate, not a process-wide global
• ADR-0008 — GPU safety: bounded waits, no early return past a barrier, runtime device limits
• ADR-0009 — A failure is an error, never a silent fallback
• ADR-0010 — The identity gate: bitwise by default, a named tolerance class otherwise
• ADR-0013 — The CPU quantizes activations to Q8_0; a device reads f32
• ADR-0014 — A KV session is a versioned, checksummed file — never a memory dump
• ADR-0021 — bf16 is a round-to-nearest-even cast, and 1-D tensors stay f32
• ADR-0030 — Launch severity lives in the helper, not in 120 call sites
• ADR-0031 — Capture runs in thread-local mode, because Global lets a foreign thread's call join the window
• ADR-0033 — The CUDA pool recycles exact byte lengths, never frees, and reports OOM as an error
• ADR-0034 — The smem opt-in is an eager pre-warm by construction, with the lazy path as defence in depth
• ADR-0035 — The q4_K dsc plane is admitted by two gates, and the payload test is equality
• ADR-0036 — bf16 weights get their own device kernels, not a dtype flag on the f16 ones
• ADR-0038 — The per-tensor registration dispatch is one shared rule, not a copy per loader