docs(path-c): sparse-MLA fused v32 e2e CORRECT on gb10 (fwd+bwd cos=1.0000, zero-copy)

The fused-fwd inf blocker is FIXED (tilelang 8385a23d: merge-pass reduce-scratch
no longer aliased onto S_shared). Re-ran scripts/verify_sparse_mla_v32_pathc_e2e.py
on gb10/sm_121a with CPPMEGA_SPARSE_MLA_V32_FUSED=1 CPPMEGA_TILELANG_CUDA_ZEROCOPY=1:

PHASE 1 (vs upstream dense ref, S=512/topk=256, ALL rows, no exclusion):
  fwd cos=1.0000 (rel 2.07e-3), dq cos=1.0000 (2.61e-3), dkv cos=1.0000 (2.35e-3)
  zero-copy bridge calls=15, eager-host=0, all finite -- real fused kernel ran, no fallback.
PHASE 2 (model scale S=4096/topk=2048): fwd+bwd finite, peak 7.45 GiB, zero-copy=15/eager=0.

path_c now has the FIRST e2e-correct fused-attention op on CUDA via DLPack zero-copy.

Perf through path_c (honest): memory win SURVIVES (12.9x less peak at S=4096/topk=256,
2.07 vs 26.8 GiB dense); the 53x LATENCY win does NOT -- per-call output torch->MLX
host bounce (_cuda_eager._torch_cuda_to_mlx .cpu().numpy(); MLX can't import a CUDA
DLPack) + per-call JIT-cache dispatch dominate wall time. Overhead-bound, not kernel-bound;
correctness is CLOSED, output-side zero-copy is the remaining pure-perf step.

Adds docs/SPARSE-MLA-PATHC-WIRED.md; updates FUSED-PIPELINE-ROADMAP.md PR-1 status.

Author: apstenku123

docs/FUSED-PIPELINE-ROADMAP.md

@@ -123,15 +123,20 @@ All paths under `/Volumes/external/sources/tilelang/examples/`. Effort tags: Low
 - **First step:** wire `examples/deepseek_v32/sparse_mla_fwd.py` + `sparse_mla_bwd.py` (+ `fp8_lighting_indexer.py` + `topk_selector.py` to feed indices) as the brick fragments for the `sparse_mla_fp8_apply{,_bwd}` descriptors in `path_c_fusion_schedules.py`, replacing the string-emitted reference. Smoke-test `atomic_addx4` (dKV) on gb10 first.
 - **Unblocks:** kills the O(n²) gather sparse-MLA on fwd AND bwd; makes v32 MLA training-capable; removes the biggest activation+compute term.
 
-**STATUS (2026-06-02, updated): WIRED + ZERO-COPY PROVEN on gb10/sm_121; the smem opt-in wall is GONE (tilelang `823c807c` re-tile fits 99 KB); e2e parity BLOCKED by a residual fused-fwd `inf` bug in the kernel itself (NOT the wiring).**
+**STATUS (2026-06-02, FINAL — e2e CORRECT): WIRED + ZERO-COPY PROVEN + e2e NUMERICALLY CORRECT on gb10/sm_121. The fused-fwd `inf` bug is FIXED (tilelang `8385a23d`); `scripts/verify_sparse_mla_v32_pathc_e2e.py` now PASSES both phases — path_c dispatches the real fused v32 sparse-MLA via DLPack zero-copy with fwd+bwd full-tensor cos=1.0000 vs the upstream dense reference, no inf/nan, no row exclusion. This is the FIRST fused-attention op proven e2e-correct in path_c on CUDA.**
 
-What changed since the previous status: the earlier "driver smem opt-in rejected ≥51200 B" wall was a codegen/argbinder smem-emission bug, **fixed** in the merge-codegen-reorg branch (`ce315509` unsized extern `__shared__` + int64 stride asserts; `67a47320`/`2524a1bb`/`823c807c` re-tile to the GB10 99 KB variant: `block_I=16`, `num_stages=1`, aggressive shared-memory merge, Hopper TMA-lower disabled, `O_shared` dropped, mask-in-shared). The fused fwd **now compiles to 93.0 KiB and LAUNCHES on real sm_121a** (CUDA 13.3); the bwd compiles + launches too. So findings (2) and (3) above are RESOLVED; finding (1) (`T.dynamic` LLVM-APInt crash) is dodged by `static_shape=True`, forced on the `gb10=True` path.
+The residual `inf` (a merge-pass reduce-scratch buffer aliased onto `S_shared`) was root-caused and fixed in tilelang `8385a23d` (Python-only change to `examples/deepseek_v32/sparse_mla_fwd.py`, no C++ rebuild). The official `test_sparse_mla_fwd`/`test_sparse_mla_bwd(check_correctness=True)` now PASS at `S=512/topk=256` and `S=4096/topk=2048` (full-tensor fwd cos 0.999998, bwd dq/dkv cos 0.999996, out_inf=out_nan=0). What changed earlier: the "driver smem opt-in rejected ≥51200 B" wall was a codegen/argbinder smem-emission bug, **fixed** in the merge-codegen-reorg branch (`ce315509` unsized extern `__shared__` + int64 stride asserts; `67a47320`/`2524a1bb`/`823c807c` re-tile to the GB10 99 KB variant: `block_I=16`, `num_stages=1`, aggressive shared-memory merge, Hopper TMA-lower disabled, `O_shared` dropped, mask-in-shared). The fused fwd **compiles to 93.0 KiB and LAUNCHES on real sm_121a** (CUDA 13.3); the bwd compiles + launches too.
+
+**E2E PARITY (`scripts/verify_sparse_mla_v32_pathc_e2e.py`, gb10/sm_121a, 2026-06-02, ALL rows, no exclusion):**
+- PHASE 1 — fused path_c vs upstream dense torch reference at `B=1 S=512 SKV=1024 H=128 HKV=1 DQK=576 DV=512 topk=256`: **fwd cos=1.0000 (rel_err 2.07e-03), dq cos=1.0000 (2.61e-03), dkv cos=1.0000 (2.35e-03)**. Zero-copy bridge calls=15, eager-host calls=0 — the real fused kernel ran, fed zero-copy, no fallback. All elements finite.
+- PHASE 2 — fused path_c at model scale `B=1 S=4096 SKV=8192 H=128 topk=2048` (no dense ref — would need ~17 TB): **fwd finite=True (3860 ms), fwd+bwd finite=True (7765 ms), peak_mem=7.45 GiB**; zero-copy=15, eager=0.
 
 - **Wiring landed (default OFF):** `cppmega_mlx/nn/_tilelang/_sparse_mla_v32_fused.py` wraps the real fused v32 `sparse_mla_fwd.py`/`sparse_mla_bwd.py` over bf16 `q/kv/indices`, wired into the differentiable `sparse_mla_fp8_path_c_apply_prepared_float` fwd + VJP, gated by `CPPMEGA_SPARSE_MLA_V32_FUSED=1`. It now (a) feeds MLX-CUDA `q/kv/indices` to the torch-backend kernel interfaces **zero-copy via DLPack** when `CPPMEGA_TILELANG_CUDA_ZEROCOPY=1` (native MLX kDLCUDA `__dlpack__` export, commit `mlx 6da6a0e4`, consumed by `torch.from_dlpack` — NO numpy host roundtrip; `_cuda_zerocopy.mlx_cuda_array_to_torch_tensor`), and (b) forwards `gb10=` (auto-detect; `CPPMEGA_SPARSE_MLA_V32_GB10` overrides) so the kernel emits the 99 KiB GB10 variant on sm_12x. RULE #1: when gated ON + `force_path_c`, every fused failure RAISES with where+what — no silent reference fallback. (Output side torch→MLX is a host bounce: MLX cannot *import* a CUDA DLPack — `convert.cpp:155` — an MLX limitation, surfaced honestly, not a degrade of the zero-copy input path.)
 - **Zero-copy: PROVEN on gb10.** `scripts/verify_sparse_mla_v32_pathc_e2e.py` traps both bridges: the fused path makes **15 zero-copy `mlx_cuda_array_to_torch_tensor` calls and 0 eager numpy-host calls** for fwd+bwd inputs. `torch.from_dlpack(mx.array)` returns a CUDA device view (`is_cuda=True`, values bit-identical) for fp32/bf16/int32. The fused kernel actually RAN (compiled `sparse_mla_fwd`/`sparse_mla_bwd_kernel`/`postprocess_kernel`, no fallback) — confirmed by the bridge counters + RULE #1 raise-on-failure.
 - **`atomic_addx4` (dKV) on sm_121: REACHED + ran** — the bwd's `sparse_mla_bwd_kernel` + `postprocess_kernel` compiled and executed on sm_121a (the dKV scatter path is no longer blocked at the smem opt-in). dKV is finite where fwd is finite.
-- **THE RESIDUAL BLOCKER — fused-fwd `inf` on a sparse subset of query rows (kernel bug, reproducible with ZERO cppmega code):** on the current gb10 tilelang HEAD (`823c807c`), `sparse_mla_fwd_interface(..., gb10=True)` produces `inf` on a **small fraction of query rows**. Pure-upstream repro (no cppmega): at `B=1,S=512,SKV=1024,H=128,DQK=576,topk=256`, the kernel emits **4864 inf elements across 19 / 65536 query rows** (LSE has 10 inf); **on the finite elements cos vs the dense torch reference = 0.9955** (i.e. the math is correct where it doesn't blow up). The example's own `test_sparse_mla_fwd(check_correctness=True)` **FAILS with `diff=nan`** at both `S=512/topk=256` and the "validated" `S=4096/topk=2048` shape — so this `inf` is present at the model shape too, and is **independent of the index distribution** (it persists even with dense-valid indices where no `block_I` tile is fully masked, and even with a guard KV row so the masked-slot gather is in-bounds). The `inf` therefore originates in the fwd online-softmax / LSE recombination (the `m_i = -(2**30)` init + rescale), not in the KV gather or the cppmega wiring. Because a single `inf` poisons any full-tensor cos/norm to `nan`, **end-to-end fwd+bwd parity-vs-reference cannot be asserted until this kernel `inf` is fixed.** Honestly reported — not faked, not silently masked.
-- **Net:** path_c's sparse-MLA op now **dispatches the real fused v32 kernel on gb10 via genuine DLPack zero-copy input** (verified: 15 zc / 0 eager, kernel ran, no fallback); the gb10 smem wall is gone. The remaining gap is a **numerical `inf` bug inside `sparse_mla_fwd` on sm_121** (reproducible via the upstream example alone) that must be root-caused before the e2e numeric proof closes. Next step: fix the fwd LSE/softmax `inf` (suspect: rows where the running-max/rescale or the `-(2**30)` sentinel overflows in bf16→fp32 accumulation on the gb10 re-tile) — a tilelang `examples/deepseek_v32/sparse_mla_fwd.py` fix, not a cppmega one. The fallback per §3.1 rank-5 (NSA) remains the alternate probe.
+- **THE RESIDUAL BLOCKER — FIXED (tilelang `8385a23d`).** The fused-fwd `inf` was a **merge-pass reduce-scratch buffer aliased onto `S_shared`** in the shared-memory merge of the GB10 re-tile: the online-softmax reduce scratch and the score buffer shared the same smem region, so a subset of query rows read corrupted partial-max state → `inf` in the LSE recombination. The fix (Python-only, in `examples/deepseek_v32/sparse_mla_fwd.py` — no C++ rebuild) keeps the reduce scratch distinct. Verified on gb10 sm_121a, ALL rows: official `test_sparse_mla_fwd`/`test_sparse_mla_bwd(check_correctness=True)` PASS; full-tensor fwd cos 0.999998, bwd dq/dkv cos 0.999996, out_inf=out_nan=0. The e2e verifier now passes (see E2E PARITY above).
+- **Latency / memory through path_c (honest, 2026-06-02):** the bare fused kernel is the standalone `53.4× faster / 26.1× less memory` win at `S=4096/topk=256` (tilelang `docs/SPARSE-MLA-GB10-BENCH.md`: fused fwd 2.86 ms / 0.259 GB vs dense 152.6 ms / 6.75 GB). **Through the current path_c wrapper the MEMORY win survives (measured 12.9× less peak at `S=4096/topk=256`: 2.07 GiB fused vs 26.8 GiB dense torch ref); the LATENCY win does NOT** — per-call wrapper overhead dominates wall time: (a) the **output** torch→MLX writeback is a host bounce (`_cuda_eager._torch_cuda_to_mlx` does `.cpu().numpy()` — MLX cannot *import* a CUDA DLPack capsule, `convert.cpp:155`), so each call copies the ~0.5 GB output GPU→host→GPU; (b) `sparse_mla_fwd_interface` does a per-call JIT-cache dispatch. Net measured fused-fwd-through-path_c ≈ 2824 ms vs dense ref 698 ms at that shape — overhead-bound, NOT kernel-bound. The **input** side is true zero-copy (0 eager / DLPack). Next perf step (separate from correctness, which is DONE): a CUDA-DLPack *import* into MLX (or an MLX-side output buffer the kernel writes in place) to kill the output host bounce, plus hoisting the kernel handle out of the per-call interface.
+- **Net:** path_c's sparse-MLA op now **dispatches the real fused v32 kernel on gb10 via genuine DLPack zero-copy input AND produces numerically-correct e2e output** (fwd+bwd full-tensor cos=1.0000 vs the upstream dense reference, 15 zc / 0 eager, kernel ran, no fallback, no row exclusion). This is the **first fused-attention op proven e2e-correct in path_c on CUDA.** Correctness is CLOSED; the remaining work is a pure-perf output-side zero-copy (above), not a correctness gap.
 
 ### PR-2 — The memory-loop fix (grad-accum + selective checkpoint barrier + 8-bit optimizer)
 **Highest mem ROI, smallest diff.** Independent of PR-1; can land in parallel.

docs/SPARSE-MLA-PATHC-WIRED.md

@@ -0,0 +1,121 @@
+# Sparse-MLA (DeepSeek-V3.2) wired into path_c — e2e CORRECT on gb10
+
+**Status (2026-06-02): DONE — e2e numerically correct.** path_c's differentiable
+sparse-MLA op dispatches the *real* fused v32 TileLang kernel (fwd **and** bwd) on
+gb10/sm_121a via genuine DLPack **zero-copy** input, and produces numerically
+correct output: full-tensor **fwd+bwd cos = 1.0000** vs the upstream dense
+reference, no inf/nan, **no row exclusion**. This is the **first fused-attention
+op proven e2e-correct in path_c on CUDA.**
+
+## What is wired
+
+- `cppmega_mlx/nn/_tilelang/_sparse_mla_v32_fused.py` — wraps the real fused v32
+  `tilelang/examples/deepseek_v32/sparse_mla_fwd.py` + `sparse_mla_bwd.py` over
+  MLX bf16 `q/kv/indices`, exposing fwd (`v32_fused_apply`) and bwd
+  (`v32_fused_bwd`, exercising the vectorized `T.atomic_addx4` dKV scatter).
+- Wired into the differentiable owner wrapper
+  `sparse_mla_fp8_path_c_apply_prepared_float` (fwd + VJP) in
+  `cppmega_mlx/nn/_tilelang/sparse_mla_fp8_path_c.py`.
+- `cppmega_mlx/nn/_tilelang/_cuda_zerocopy.py` — native MLX kDLCUDA `__dlpack__`
+  export → `torch.from_dlpack`, NO numpy host roundtrip. The torch tensor is a
+  real CUDA device view of the MLX allocation.
+
+### Gates (both required for the fused zero-copy path)
+
+```
+CPPMEGA_SPARSE_MLA_V32_FUSED=1     # dispatch the real fused v32 kernel
+CPPMEGA_TILELANG_CUDA_ZEROCOPY=1   # feed inputs zero-copy via DLPack
+# CPPMEGA_SPARSE_MLA_V32_GB10=1    # optional: force the 99 KiB GB10 re-tile
+```
+
+Default (gates OFF): callers keep the existing reference path; nothing fused runs.
+
+**RULE #1 (no silent fallback):** when gated ON + `force_path_c`, every fused
+failure RAISES with where+what. There is no silent reference fallback.
+
+## Layout / ABI
+
+- `q`  : bf16 `(B, S, H, DQK)`, `DQK = d_v + tail = 512 + 64 = 576`.
+- `kv` : bf16 `(B, SKV, G, DQK)` (kv_group `G`).
+- `indices` : int32 `(B, S, G, topk)` (sentinel `SKV` → masked slot).
+- fwd → `out` bf16 `(B, S, H, d_v)`, `lse` fp32 `(B, S, H)`.
+- bwd → `dq` bf16 `(B, S, H, DQK)`, `dkv` bf16 `(B, SKV, G, DQK)`.
+
+## E2E verification
+
+`scripts/verify_sparse_mla_v32_pathc_e2e.py` — traps BOTH input bridges
+(zero-copy and eager-host) so it can prove which one fed the kernel, runs the
+path_c owner wrapper through fwd + `mx.grad` VJP, and asserts parity + finiteness
+over the **full tensor** (RULE #1: any non-finite element raises with the exact
+offending query rows; any eager-bridge call raises).
+
+Run on gb10:
+
+```
+cd /home/dave/source/cppmega_mlx
+CPPMEGA_SPARSE_MLA_V32_FUSED=1 CPPMEGA_TILELANG_CUDA_ZEROCOPY=1 \
+  /home/dave/cppmega-venv/bin/python scripts/verify_sparse_mla_v32_pathc_e2e.py
+```
+
+### Results (gb10 / sm_121a, tilelang `8385a23d`, 2026-06-02, ALL rows)
+
+**PHASE 1 — fused path_c vs upstream dense torch reference**
+`B=1 S=512 SKV=1024 H=128 HKV=1 DQK=576 DV=512 topk=256`:
+
+| tensor | cos | rel_err |
+|--------|-----|---------|
+| fwd | 1.0000 | 2.07e-03 |
+| dq  | 1.0000 | 2.61e-03 |
+| dkv | 1.0000 | 2.35e-03 |
+
+zero-copy bridge calls = **15**, eager-host bridge calls = **0**, all elements finite.
+
+**PHASE 2 — fused path_c at model scale** (no dense ref — it would need ~17 TB)
+`B=1 S=4096 SKV=8192 H=128 topk=2048`:
+
+- fwd: finite=True, latency = 3860 ms, out_abs_mean = 0.0140
+- fwd+bwd: finite=True, latency = 7765 ms, dq_abs_mean = 0.0051, dkv_abs_mean = 0.0571
+- peak_mem = **7.45 GiB**
+- zero-copy = 15, eager = 0
+
+The proof that the **real fused kernel ran** (not the reference, not the eager
+host bridge): `zero-copy bridge calls = 15` and `eager = 0` for fwd+bwd inputs,
+plus the TileLang compile logs (`sparse_mla_fwd` / `sparse_mla_bwd_kernel` /
+`preprocess_kernel` / `postprocess_kernel`) — and RULE #1 raise-on-failure means
+a fallback could not have masked anything.
+
+## The fixed `inf` bug (tilelang `8385a23d`)
+
+The earlier blocker — `inf` on a sparse subset of query rows in the GB10 re-tile
+fwd — was a **merge-pass reduce-scratch buffer aliased onto `S_shared`**: the
+online-softmax reduce scratch and the score buffer shared the same shared-memory
+region, corrupting partial-max state for some rows → `inf` in the LSE
+recombination. Fixed Python-only in `examples/deepseek_v32/sparse_mla_fwd.py`
+(no C++ rebuild). Official `test_sparse_mla_fwd`/`test_sparse_mla_bwd
+(check_correctness=True)` now PASS at `S=512/topk=256` and `S=4096/topk=2048`;
+full-tensor fwd cos 0.999998, bwd dq/dkv cos 0.999996, out_inf = out_nan = 0.
+
+## Performance through path_c (honest)
+
+The **bare** fused kernel is the standalone `53.4× faster / 26.1× less memory`
+win (tilelang `docs/SPARSE-MLA-GB10-BENCH.md`, `S=4096/topk=256`: fused fwd
+2.86 ms / 0.259 GB vs dense 152.6 ms / 6.75 GB).
+
+Through the **current path_c wrapper**:
+
+- **Memory win SURVIVES** — measured **12.9× less peak** at `S=4096/topk=256`
+  (2.07 GiB fused vs 26.8 GiB dense torch reference).
+- **Latency win does NOT survive** — per-call wrapper overhead dominates:
+  - the **output** torch→MLX writeback is a host bounce
+    (`_cuda_eager._torch_cuda_to_mlx` does `.cpu().numpy()` — MLX cannot *import*
+    a CUDA DLPack capsule, `convert.cpp:155`), copying the ~0.5 GB output
+    GPU→host→GPU every call;
+  - `sparse_mla_fwd_interface` does a per-call JIT-cache dispatch.
+  - Net measured fused-fwd-through-path_c ≈ 2824 ms vs dense ref 698 ms at that
+    shape — **overhead-bound, NOT kernel-bound.** The **input** side is true
+    zero-copy (0 eager / DLPack).
+
+**Correctness is CLOSED.** The remaining work is pure-perf: a CUDA-DLPack
+*import* into MLX (or an MLX-side output buffer the kernel writes in place) to
+kill the output host bounce, plus hoisting the kernel handle out of the per-call
+interface. That is a latency optimization, not a correctness gap.