feat: shared-engine refactor + Nemotron-Nano-30B GB10 prefill/decode optimizations (~5×)

[TheTom]

First time here? Read the TL;DR + "where to start" map below — you don't need to read all 104 commits. Big draft (159 files, ~29K lines) opened for visibility.

TL;DR

Two things, both backend-gated so the default path can't regress:

1. Shared-engine refactor — each model's forward is written once and runs on both CUDA (server GPU) and Metal (Apple) backends; +ffai-vulkan for AMD.
2. A deep prefill/decode optimization pass on Nemotron-Nano-30B (Mamba2 + 128-expert-MoE hybrid) on a single GB10.

Prefill result so far (S=2048, GB10, byte-matched precision, argmax-gated)

context depth	start of pass	now
d0	187 tok/s	~411–446
d32768	103 (collapsing)	~334 (flat)

The curve went from collapsing to flat ~370 across d0→d32768 — ~2.3× at d0, ~3.4× at deep context. Every step correctness-gated (logit cosine + shared-top-5, not just argmax).

The levers (each gated, each measured)

• Conv double-shift fix — a real correctness bug: the all-device Mamba decode step shifted the causal-conv ring twice, corrupting multi-token decode past token 0 (invisible at pos 0, so Phase-1 missed it).
• SSD matmul scan (NEMOTRON_SSD_MATMUL) — Mamba2 state-space-duality chunked-matmul via batched cuBLAS; ssm_scan stage 26%→5% (6.8×), +36% e2e @ d0. Scan-level cosine 1.000000.
• Tensor-core flash-attention (sdpa_multi_tc, depth auto-select) — cuBLAS online-softmax, GQA stride-0 fan-out; sdpa stage 86%→30% @ d32768 (up to 14×), +59% deep-context e2e.
• Launch-geometry fix — six prefill elementwise kernels were dispatched block:[1,1,1] (one thread/block!); packed into real blocks → +7–13%, byte-identical.
• Batched prefill on Metal (backend-gated) — first Nemotron-30B prefill on Apple Silicon (M5 Max, 33.9 tok/s @ S=2048), resident weights, EXACT-MATCH vs reference. CUDA path untouched. Adds an MLX-4bit loader.
• Earlier: cuBLAS tensor-core GEMM escape hatch, fully on-device MoE, batched rope/kv/conv, W4A16 MoE GEMM family, deep-context KV-cap IMA fix.

Honestly scoped

• The MoE expert GEMM is settled at per-expert cuBLAS (~27 TFLOP/s) — three faster-kernel attempts (TC-grouped, on-device, grouped-Marlin) all proven slower; it's memory-pipe-bound, would need a multi-day CUTLASS-grade rewrite. Profiling shows the next real win is the GPU-idle gap (per-forward cuMemAlloc/cuMemFree device-syncs → a caching allocator, in progress), not the GEMM.
• vs a mature serving engine's prefill (~6,400 t/s) we're not there yet — but the curve is flat and every gain is measurement-backed.

Where to start reviewing

• Refactor: ffai-modeltests/ (forward written once), ffai-runtime/, ffai-loader/ (+MLX dequant).
• Prefill levers: ffai-ops/src/lib.rs (kernels: sdpa_multi_tc, ssm_prefill_scan_ssd, launch-geometry) + ffai-modeltests/src/lib.rs forward_batched (gates + backend select).
• Big line-count is per-backend test/bench files — skim.

Build/validate

cargo test --release -p ffai-cuda --features cuda --test nemotron; per-op profiler NEMOTRON_PROFILE=1; correctness NEMOTRON_PREFILL_CHECK=1. Pairs with the metaltile kernel PR. WIP — happy to walk through live.

[TheTom]

🏆 d0 prefill target BROKEN: 6534-6579 tok/s sustained vs the 6395 published number (102-103%)

Five consecutive runs ≥ 6534, argmax fingerprint stable (1294), same-window controls, co-resident server idle. The full ladder from this morning's 1912:

lever	d0 tok/s
day start (bit-exact baseline)	1912
native FP4 W4A4 grouped MoE + day's kernel work	2868
host-bubble + upload fixes	4444
cp.async mma + CUDA-graph replay (bit-exact, oracles green)	4778
CUTLASS grouped MoE (device descriptors)	4994
+ FP8/FP4 dense projections	5331
+ Lt FP4 shared expert	6173
+ graph compose + f32-direct quant	6305
+ coalesced transposed SSD recurrence	6545-6579

= 3.4× in one day, with the bit-exact path preserved separately at 4778 (every depth oracle green).

Depth rows after split-KV attention (f1d81d6): d4096 89% / d8192 80% / d16384 74% / d32768 70% of their targets. Measured analysis: the f16 attention mma already runs near the chip's f16 tensor-core peak at depth — the remaining depth lever is an FP8 attention path (2× mma rate + FP8 KV cache, which is the precision this model family's published recipe ships anyway). That's the next project.

[TheTom]

Depth attention diagnosis corrected (no perf change committed — recording the finding). Profiled d32768: attention is 82% of the forward and runs at only ~8 TFLOP/s (3-5% MFU) — so it is occupancy/latency-bound, not compute-bound (my earlier 'near mma peak' read was wrong). Confirmed by experiment: widening the KV tile BC 16→32 made depth slower (more smem/block → fewer blocks/SM → the serial KV walk's latency stops hiding), so split-KV / fp8-mma / cp.async all measured par for the same reason — none raised occupancy. Reverted the BC=32 try; tree stays at the committed cp.async kernel.

The real depth lever is a higher-occupancy attention kernel (fewer smem bytes/block, or a warp-specialized KV-producer / mma-consumer FMHA) — a named follow-up project, ~2-4× attention at depth on paper, which would cross all the depth rows. d0 is unaffected and remains broken at 6534-6579.

[TheTom]

⚠️ Correction — quality gate retracts the 6579 headline

Ran the logit-agreement gate (cosine + top-5 overlap vs the sequential reference) on the configs, which I should have done before quoting the stacked numbers. Result:

config	d0 tok/s	argmax	cosine	verdict
CUTLASS-MoE + graph + recur_t	~5190	1104 ✓	0.971	quality-preserving
+ FP8 mamba	~5300	1886 ✗	0.945	argmax flips
full stack (+FP4-dense=all)	6579	1294	0.44	output corrupt (0/5 top-5)

The 6579 rode on FP4-quantizing the dense/shared projections, which collapses logit cosine to 0.44 — the 'stable 1294 argmax' was deterministic garbage, not preserved quality. Retracting it as a headline.

Honest standing: the quality-preserving (argmax-exact, native-FP4-MoE) best is ~5190 t/s at d0 = 81% of the 6395 target. That's the real, shippable number; the FP4/FP8-dense levers stay opt-in-and-flagged-unsafe. recur_t and the CUTLASS device-descriptor MoE are both quality-safe and carry the win. Re-aiming the remaining work at quality-neutral levers only (the MoE GEMM at 77ms and the SSD elementwise chain at ~92ms dominate the forward now).

[TheTom]

Quality-fix investigation: dense-projection quantization can't be made fast+quality on this part (closing the avenue honestly).

Chased the proper fix for the dense-projection speed lever:

• per-tensor FP8 (mamba in/out): cosine 0.945, argmax flips — too coarse for the 2688-wide projections (outliers). (It was never clean; earlier 'argmax-exact' only checked argmax, not the logit distribution.)
• per-channel FP8 (per-output-channel weight + per-token act, the industry-standard W8A8): cuBLASLt returns NOT_SUPPORTED for OUTER_VEC FP8 scaling on sm_121 — the library doesn't expose per-channel FP8 GEMM on this GPU.
• FP4 dense: cosine 0.44 (ue4m3 subnormal-floor + FP4 is simply too aggressive for these layers regardless of scale fixes).

Conclusion: the only quantization this model tolerates with quality intact is the FP4 MoE (already shipped, argmax-exact to the HF oracle at every depth, cosine 0.971) + the CUTLASS device-descriptor path. Dense/shared/attention projections must stay f16 — which matches NVIDIA's own recipe keeping the attention-adjacent layers at BF16.

So the quality-preserving prefill ceiling via quantization is ~5200 t/s d0 (81% of 6395), and that's honest-final. Pushing past it with quality intact is now purely kernel-structural: the SSD elementwise-fusion chain (fold the decay·dt scaling passes into adjacent GEMM epilogues) and a higher-occupancy FMHA for the depth rows. Reverted the non-working per-channel code; tree is clean at the quality-safe config.

[TheTom]

Reference model findings — NVFP4-QAT checkpoint (for posterity)

Pulled and analyzed the published reference checkpoint behind the arena pp2048 ladder we're chasing (nvidia/Nemotron-3-Nano-Omni-30B-A3B-Reasoning-NVFP4, 22.4 GB, 3 shards). Recording the precision recipe and how the reference number was produced — it explains both our earlier quality wall and the remaining speed gap.

Reference throughput (pp2048, c1): d0 6395.5, @D4096 5482.6, @d8192 4993.3, @d16384 3979.9, @d32768 2734.5; tg128 ≈ 57 (we already beat decode). Served with chunked prefill at max_num_batched_tokens=32768.

Checkpoint precision recipe (MIXED_PRECISION, modelopt):

• MoE expert up_proj/down_proj — NVFP4: weight U8 packed e2m1 [N, K/2], weight_scale E4M3 block-16 ue4m3 [N, K/16], weight_scale_2 F32 per-tensor global, input_scale F32.
• Mamba in_proj/out_proj — FP8 e4m3: weight E4M3 + weight_scale F32 (per-tensor) + input_scale F32.
• KV cache — FP8.
• conv1d.weight, A_log, norms — BF16 (kept high-precision).

Why this matters for us:

1. The two-level MoE scale (block ue4m3 plus per-tensor global weight_scale_2) is exactly what our earlier post-hoc BF16→FP4 dense quant lacked — the missing global is what drove the ue4m3 subnormal-floor bug (cosine 0.44 → 0.84). Our current MoE path quantizes with a proper two-level scale and is already quality-clean.
2. The weights are QAT-trained for FP4/FP8, so the reference runs mamba in_proj/out_proj in FP8 with no quality loss. Our post-hoc FP8-dense was either lossy (cosine 0.945, per-tensor) or net-slower (per-channel: argmax-exact but quant + post-scale overhead exceeded the FP8 saving). QAT weights remove the runtime weight-quant entirely → FP8 mamba becomes both clean and a net win.

Current quality-clean baseline (this branch, GB10, CUDA-graph replay): pp2048 d0 = 5206 tok/s, argmax-exact, cosine 0.9714 = 81% of the reference 6395.

Per-op profile @ S=2048 (sync-bracketed, so absolute ms inflated — proportions are real):

op	%	note
moe_experts	25.7%	NVFP4 grouped GEMM, 48.5 TFLOP/s (4.85% peak)
proj_gemm	24.0%	dense f16 (mamba + attn proj) — the FP8 lever lives here
ssm_scan	15.9%	SSD chunked-matmul chain, 1.4 TFLOP/s — bandwidth/launch-bound
slice/cast	13.9%	208 f32↔f16 cast launches — pure pass-hygiene overhead
moe_shared	5.9%	shared expert, 70.5 TFLOP/s
sdpa_prefill	5.9%	attention
conv/gated_norm/rms	~7%	elementwise

The gap to 6395 is distributed, not a single kernel: MoE GEMM tiling + FP8 mamba proj (needs the QAT weights) + SSD elementwise fusion + collapsing the 208 slice/cast launches. Next steps: wire native FP8 mamba in/out_proj from the QAT checkpoint (clean + ~2× on the 24% proj_gemm slice), and fuse away the cast launches.

[TheTom]

357ddb2 — Marlin W4A16 MoE correctness (fully on-device PATH B)

This makes the NEMOTRON_W4A16_MARLIN routed-expert path produce correct, quality-safe output with no host round-trips (CUDA-graph-capturable), layered on top of the existing prefill stack.

Root cause fixed. Under the device-descriptor routing path the host triples vector is empty (routing offsets/sorted-tokens/weights live on-device in moe_dev), but the Marlin path built its per-row expert ids, gather, and scatter from triples — so idx_dev was a NULL device pointer and moe_w4a16_marlin did an illegal read at 0x0 (localized with compute-sanitizer + a device-pointer dump).

Changes

• moe_expert_ids_from_offsets: new device kernel expanding CSR offsets [n_exp+1] into per-row expert ids [mt] — graph-safe replacement for the host build.
• on-device PATH B: gather via the router sorted_tok (st_dev), scatter via moe_scatter_add_det_dev, accumulator kept on device and consumed by the device shared-expert merge (acc kept on device since ondevice_moe is hardwired off for use_w4a16).
• fp4_skip_gather gated with !use_w4a16 so Marlin keeps its own gather.
• tile-select fix: the 128-tile is used only when both n_out and k_in are multiples of 128 — the down-GEMM k_in=1856 must use the 64-tile, else the 128-strided K-loop reads past the weight buffer.
• relu2_scale_f16: clamp pre-square to the f16-safe range (no Inf/NaN).
• dtype-aware dl: cast F16 tensors to F32 before host download.

Measured (GB10, pp2048, NemotronH-Nano-30B-A3B). Batched argmax matches the sequential reference exactly (1104), logit cosine 0.962, no NaN. This is the quality-safe direction — weight-only W4A16 preserves activation precision, vs FP4-dense at cosine 0.84 and the W4A4 LT shared expert. Quality verified eager; graphed throughput (combining with NEMOTRON_W4A16_SHARED for an all-W4A16 graph-safe path + the SSD-recurrence speed stack) is the next measurement.

[TheTom]

bd34001 — keep the on-device Marlin forward CUDA-graph-capturable

Follow-up to 357ddb2. The xn_h_opt host download (dl(&xn)) is computed whenever use_w4a16 is set (to prep the host gather), but the on-device Marlin path (PATH B) gathers via the router sorted_tok and never reads xn_h — yet that dl still fired and aborts CUDA-graph stream capture. Guarded it with !(use_w4a16_marlin && moe_dev.is_some()).

Verified on GB10: with this guard the Marlin + NEMOTRON_W4A16_SHARED forward captures into a CUDA graph and replays argmax-exact (1104). Throughput of the W4A16 path is still GEMM-bound (the 64-tile WMMA runs at ~0.66% of peak because it reads scales per-nibble; the scale-preloaded 128-tile needs inter padded 1856→1920 to engage — that speedup is in progress). Landing the capture fix now since it's independent and correct.

[TheTom]

Progress update — FP4-dense NaN root-caused + fixed (runs at 6359 tok/s); quality work in progress

Since 357ddb2/bd34001 (Marlin W4A16, cosine 0.962 argmax-exact + graph-capture fix), the focus has been pushing prefill higher while holding quality. Findings on GB10/SM121 @ pp2048:

FP4-dense NaN root-caused to a documented upstream bug. The cutlass FP4 dense GEMM produces NaN/garbage on GB10 for the large projections (in/out/o; q/k/v with K=2688 survive). Root cause = the SM120/121 FP4 tile configs need >99 KiB SMEM but GB10 has only 99 KiB → SMEM overflow (NVIDIA/TensorRT-LLM#11368; corroborated by cutlass#3096, flashinfer#2723/#2577, sgl-project/sglang#18954). This is exactly why the vendor recipe FP8s in/out/o.

Fix + speed: routed dense FP4 through our own GB10-sized hand grouped-MMA kernel (the one the MoE FP4 path already uses, n_exp=1) to dodge the SMEM overflow. Result: FP4-dense runs on GB10 at ~6359 tok/s graphed (vs FP8-dense 5853, +8.6%), no NaN.

Quality at 6359 = cosine 0.90 (not yet shippable): the in/out/o activations carry systematic per-channel outliers that FP4 loses. Per-token activation scales are the wrong fix (per-token ≠ per-channel; measured a regression). Pursuing a Hadamard rotation (QuaRot/SpinQuant-style) to spread the per-channel outliers — first prototype regressed (likely the non-zero-mean DC-concentration issue), debugging the exactness now.

Status: Marlin W4A16 (357ddb2) remains the argmax-exact quality path (5853-class). FP4-dense-hand (6359) is a real speed lever + a fix to a documented cutlass-SM121 bug, gated pending the outlier-quality work. Will push the dense-hand path once the quality lands (target ~6359 @ ≥0.95). Code staged; nothing half-baked pushed.

[TheTom]

Catch-up commits with the ffai side of the GB10 Nemotron-Nano-30B prefill effort:

• 0324c23 feat(nemotron): W4A8/W8A8 MoE paths, Q4-W4A16 kernel speedups, SSD scan.
  • W4A8 and W8A8 grouped MoE paths: ffai-core trait methods, ffai-cuda backend impls, ffai-ops wrappers, and the NEMOTRON_W4A8_MOE / NEMOTRON_W8A8_MOE gates. W8A8 is argmax-exact on the 30B prefill with near-lossless fp8 weights.
  • Q4-native W4A16 MoE kernel (moe_q4_grouped_mma): BN=128 tiling to cut redundant activation reads, plus A-operand ldmatrix. About +52 percent over the prior W4A16 path at equal quality.
  • Mamba2 SSD scan kernels with a correctness test, and standalone W4A8/W8A8 grouped-GEMM tests.
• 41be7bc test: fp8 projection-GEMM microbench plus f16-norm, slice, and qwen MoE test files.

The CUDA path depends on the matching metaltile FFI on thewafflehaus/iron#13 (branch tom/cuda/nvfp4-moe-gemm); the metaltile git dependency should point there to compile the W4A8/W8A8 calls.

[TheTom]

NVFP4 grouped-MoE GEMM: standalone correctness + speed validated on GB10 (sm_121a)

Validated the persistent-grid NVFP4 grouped-MoE GEMM end to end at the Nemotron-Nano-30B MoE shape (M=2048, K=2688, N=1856, 128 experts, top-6, group-16):

• Correctness: cosine 0.999994 vs an f32 dequant reference (output 22,806,527 / 22,806,528 nonzero, matching norm).
• Throughput: 3.047 ms/iter, 40.23 TFLOP/s, which is 3.12x the current hand-rolled W4A16 grouped-MMA kernel (~9.51 ms) at near-bf16 quality.

Root cause of an earlier all-zero output (now resolved): the 4-bit dequant emits deliberately bias-shifted small values, so the per-group fp8 scale must be stored in a bias-compensating layout (permute, multiply by a shared power-of-two scale_factor, multiply by 2^7, reinterpret as int16 and shift left 1, take the high fp8 byte), with the per-tensor global scale folded through the top-k weight path and divided by scale_factor. With nominal fp8 scales the scale-times-weight product underflows to zero. Diagnosed by sentinel-prefill, force-nonzero-operands, and an MMA fragment trace.

Builds for sm_121a via relocatable device code (whole-program device link).

Next: wire the weight repack, the scale prep, and the grouped GEMM call into the MoE model path here, replacing the hand kernel, then measure 30B end to end. Amdahl puts the MoE-GEMM swap alone at roughly 3300 to 3630 tok/s; reaching the prefill target needs it stacked with CUDA-graph capture plus the cast, router, and norm fusions (tracked separately).

[TheTom]

Milestone on the NVFP4 MoE GEMM effort: the F16 path is validated at the Nemotron MoE shape with cosine 1.000000 vs an f32 reference and 40 TFLOP/s (3.12x the hand kernel), and the host weight-prep (quantize, transpose, S0E5M3 scale transform, global) is byte-identical to the reference in a Rust test.

The ffai-side wiring (core trait + cuda impl + ops wrappers + the S0E5M3 transform) is in place and compiling against the device static lib. The model gate (routing + prepare loop + forward) plus the 30B end-to-end are the remaining steps; the ffai-side code lands here as the gate completes. Full writeup + the kernel side are on the metaltile PR (research/nvfp4-moe-gemm/FINDINGS.md).

[TheTom]

Progress update: Nemotron-Nano-30B GB10 prefill (S=2048)

Current best correct config: 4672 tok/s at logit cosine 0.97138 vs the bf16 reference, last-token argmax exact (1104). About 4870 tok/s without the per-step correctness gate. Reference target on the same GB10: 6395 tok/s, so the remaining gap is about 37 percent.

Recent fixes and wins:

• Fixed an illegal-memory-access crash in the block-scaled NVFP4 grouped-GEMM prefill path. The gathered activation buffer was being skipped for that path (a guard meant for the gather-on-read variant did not exclude it), so the grouped GEMM read a 1-element dummy tensor and ran off the end. Standalone unit test at the real expert shape passes at max-rel 3.4e-4; the fix restores the full path to argmax-exact.
• Mamba2 SSD chunk length: 128 is the sweet spot (64 gives 4591, 128 gives 4672, 256 gives 4594). Both speed and cosine improved at 128.

Profile at S=2048 (relative): dense projections 22 percent, routed experts 22 percent, SSD scan 18 percent, split and cast 14 percent, shared expert 8 percent, attention 7 percent.

Negative results (logged so we do not repeat):

• Per-tensor fp8 dense projections: the GEMM is 1.5 to 1.9x faster in isolation at cosine 0.999, but on the full model the per-tensor activation scale flips the argmax (1104 to 1350) because this model has heavy activation outliers. Per-channel holds the argmax but the post-pass eats the speedup and cosine drops to 0.938. The projection weights are also Q4-sourced here, so they double-quantize.
• Swapping the routed-expert GEMM kernel does not move end-to-end throughput (all variants cluster near 4400 to 4672), so the experts are not the prefill bottleneck.

Next levers (each a real build): activation-outlier migration to enable fp8 activations at quality (the single biggest lever, unlocks both projections and experts), and op-fusion to cut the launch count (conv plus split plus scan, q/k/v concat, fused epilogues).

[TheTom]

Progress update: Nemotron-Nano-30B GB10 prefill (S=2048), continued

New best correct config: about 5130 tok/s (5088 to 5156 across runs, 5099 on a fresh cool measure) at logit cosine 0.952, last-token argmax exact (1104, with a top1 minus top2 margin of 1.26). That is +10 percent over the 4672 reported above, and it stacks cleanly from four levers:

• Fused residual-add plus next-layer RMSNorm in one pass: 4672 to 4890.
• Per-tensor fp8 dense projections held to argmax by a block Walsh-Hadamard rotation applied to the activation before quantization: 4890 to 5085. Plain per-tensor fp8 flips the argmax on this outlier-heavy model; the rotation makes per-tensor hold (cosine 0.952). A hard activation clip is about 2 percent faster but its argmax is non-deterministic run to run, so the rotation is the trustworthy path.
• Fused intra-chunk SSD plus shared-expert MMA: 5085 to 5130.

Single-shot prefill on this GB10 is about plus or minus 10 percent run to run (power-capped), so every comparison here is cool-run vs cool-run.

In progress, not yet a measured win: running the residual and hidden stream in 16-bit instead of f32, to halve the memory traffic through every norm, add, and elementwise op. The full boundary cascade is now wired and runs end to end with no faults. A standalone test confirms the 16-bit norm and the 16-bit tensor-core GEMM are numerically clean (cosine 0.999996 and 0.999999 vs f32 at the model shapes), so the kernels are sound. The remaining work is a set of dtype boundary casts in the MoE block where several call sites still assume an f32 activation. No throughput number on this path yet; will report once it is argmax-correct.

All of the above is flag-gated. The default path is unchanged and re-verified at 5099 tok/s, cosine 0.952, argmax 1104, so there is no regression risk from the staged work.

[TheTom]

Closing as superseded: this combined backends branch has been fully decomposed into focused, individually CI-green PRs.

Where everything landed:

• Rust inference engine (rust/, 12 crates + metal/cuda/vulkan backends) → feat(rust): add the Rust inference engine under rust/ #22 (merged), built against the org metaltile dev git dep with CUDA/Vulkan feature-gated off and a path-filtered rust.yml.
• Kernel-reorg call-site migration (the mt_* rename, rope_banded / partial_rope / grouped-SSM contracts) → fix: sync FFAI to the metaltile kernel reorg (binding-gen + call-site migration) #21 (merged).
• DSv4 / GGUF / prefill / AURA / telemetry Swift source → already merged via feat(aura): TQ+ port — KLD harness + 2-pass compressed flash + unified-dtype cache #15, feat(step3): Step-3.5/3.7-Flash family scaffold + config decoders #16, feat(gguf,dsv4): GGUF v3 reader + DeepSeek-V4-Flash GGUF loader & forward path #17, feat(telemetry): consolidate quality metrics behind one flag-gated tap #19, fix: sync FFAI to the metaltile kernel reorg (binding-gen + call-site migration) #21. (Stats/Perplexity.swift here is the pre-relocation duplicate of the merged Telemetry/Perplexity.swift.)
• gemv per-dispatch overhead micro-bench → test: add gemv per-dispatch overhead micro-bench #23 (merged), the last genuinely-new functional file.

A prerequisite metaltile fix surfaced while validating the Rust port (three q4 coalesced gemv kernels typed their scale f16 instead of f32) landed as thewafflehaus/iron#33.

Intentionally not carried over: the architecture/perf/verification docs and the Nemotron research/profiling notes (kept out for now), and IndirectDispatchTests (targets a kernel variant removed in the reorg).