feat(path-c): wire fused v32 sparse-MLA to zero-copy DLPack on gb10 (kernel inf blocks e2e)

apstenku123 · apstenku123 · commit 1b021e86434e · 2026-06-02T19:46:52.000+02:00
Wire `_sparse_mla_v32_fused.py` (already env-gated into the differentiable `sparse_mla_fp8_path_c_apply_prepared_float` fwd+VJP) to feed MLX-CUDA q/kv/indices to the torch-backend v32 kernel interfaces ZERO-COPY via DLPack, and forward the `gb10` arch flag so the kernel emits the 99 KiB sm_121 variant. Wiring (correct + verified): - `_cuda_zerocopy.mlx_cuda_array_to_torch_tensor`: MLX-CUDA -> torch CUDA tensor zero-copy. Prefers the native MLX kDLCUDA `__dlpack__` export (mlx 6da6a0e4, `torch.from_dlpack(mx.array)`); falls back to the `data_ptr()` capsule builder only when native export is absent. Both are genuine device-view zero-copy (NOT a silent host-copy degrade). RAISES on failure (RULE #1). - `_sparse_mla_v32_fused._mlx_to_torch_cuda`: under CPPMEGA_TILELANG_CUDA_ZEROCOPY=1 uses the zero-copy bridge instead of the numpy-host-roundtrip eager bridge. - forward `gb10=` to `sparse_mla_fwd_interface` / `sparse_mla_bwd` (auto-detect by default; `CPPMEGA_SPARSE_MLA_V32_GB10` overrides). - fix a torch-tensor bug: dkv cast used MLX/numpy `.astype` -> `.to(torch.bfloat16)`. - scripts/verify_sparse_mla_v32_pathc_e2e.py: traps both bridges to PROVE the fused kernel ran via zero-copy (15 zc-calls / 0 eager), surfaces a kernel inf blowup loudly with the exact non-finite rows. PROVEN on gb10 (sm_121a, CUDA 13.3, tilelang 823c807c): - zero-copy input bridge: 15 `mlx_cuda_array_to_torch_tensor` calls, 0 eager numpy-host calls, for the fwd+bwd path_c op. `torch.from_dlpack(mx.array)` returns a CUDA view (is_cuda, bit-identical) for fp32/bf16/int32. - the real fused fwd+bwd kernels compiled + LAUNCHED on sm_121a (the prior >48 KB dynamic-smem opt-in wall is GONE after the 823c807c re-tile; 93.0 KiB fwd fit). atomic_addx4 dKV scatter reached + ran. BLOCKER (kernel, not wiring): `sparse_mla_fwd_interface(..., gb10=True)` on this gb10 HEAD produces `inf` on a sparse subset of query rows (e.g. 4864 inf over 19/65536 rows at S=512/topk=256; cos=0.9955 on the FINITE elements). Reproducible with ZERO cppmega code via the upstream example's own `test_sparse_mla_fwd` (`diff=nan`) at both S=512/topk=256 and the model S=4096/topk=2048 shape; persists with dense-valid indices and a guard KV row, so it is a fwd LSE/online-softmax overflow inside the kernel, not the gather/mask or the wiring. A single inf poisons full-tensor cos -> nan, so e2e fwd+bwd parity-vs-reference cannot be asserted until the tilelang `sparse_mla_fwd.py` inf is fixed. Reported honestly, no silent reference fallback (RULE #1).
diff --git a/cppmega_mlx/nn/_tilelang/_cuda_zerocopy.py b/cppmega_mlx/nn/_tilelang/_cuda_zerocopy.py
@@ -287,3 +287,90 @@ def mlx_cuda_array_to_tvm_tensor(arr: "mx.array") -> Any:
             f"dtype={arr.dtype}, shape={tuple(int(d) for d in arr.shape)}): "
             f"{type(exc).__name__}: {exc}"
         ) from exc
+
+
+def _mlx_native_cuda_dlpack_available(arr: "mx.array") -> bool:
+    """Return True iff this MLX build natively exports a kDLCUDA DLPack capsule.
+
+    The native C++ export (MLX ``feat(dlpack)`` / commit 6da6a0e4) advertises
+    ``__dlpack_device__() == (kDLCUDA(2), id)`` and produces a real CUDA capsule
+    from ``__dlpack__()``. Older MLX builds either lack ``__dlpack__`` or raise
+    "CUDA DLPack export is not supported", in which case we fall through to the
+    ``data_ptr()`` Python escape hatch below. Both are genuine zero-copy paths.
+    """
+
+    dl_dev = getattr(arr, "__dlpack_device__", None)
+    if not callable(dl_dev) or not callable(getattr(arr, "__dlpack__", None)):
+        return False
+    try:
+        dev = dl_dev()
+    except Exception:
+        return False
+    return bool(dev) and int(dev[0]) in (kDLCUDA, kDLCUDAManaged)
+
+
+def mlx_cuda_array_to_torch_tensor(arr: "mx.array") -> Any:
+    """Import an MLX-CUDA array as a torch CUDA tensor view, zero-copy, via DLPack.
+
+    This is the device-view (no host roundtrip) counterpart to the numpy-copy
+    bridge in ``_cuda_eager._mlx_to_torch_cuda``. It feeds the TileLang
+    torch-backend kernel interfaces (``sparse_mla_fwd_interface`` /
+    ``sparse_mla_bwd``) a real device view of the MLX allocation — no copy, no
+    host bounce.
+
+    Two genuine zero-copy mechanisms, in preference order (NOT a silent
+    degrade — both are real zero-copy device views):
+
+    1. **Native MLX kDLCUDA export** (commit 6da6a0e4): hand the MLX array
+       straight to ``torch.from_dlpack``, which calls its ``__dlpack__()``.
+    2. **``data_ptr()`` Python escape hatch**: build the kDLCUDA
+       ``DLManagedTensor`` capsule from ``mx.array.data_ptr()`` (MLX #3342) and
+       import it. Used only when the native export is absent.
+
+    RAISES (no copy fallback) so a broken zero-copy path is surfaced, per RULE #1.
+    """
+
+    import torch
+    from torch.utils.dlpack import from_dlpack as _torch_from_dlpack
+
+    if _mlx_native_cuda_dlpack_available(arr):
+        # Native zero-copy: torch.from_dlpack consumes MLX's own __dlpack__()
+        # kDLCUDA capsule. MLX must flush its CUDA stream before handoff.
+        _synchronize_mlx_stream()
+        try:
+            t = torch.from_dlpack(arr)
+        except Exception as exc:  # noqa: BLE001 - surface where + what failed
+            raise RuntimeError(
+                f"_cuda_zerocopy: torch.from_dlpack rejected the MLX native "
+                f"kDLCUDA export (dtype={arr.dtype}, "
+                f"shape={tuple(int(d) for d in arr.shape)}): "
+                f"{type(exc).__name__}: {exc}"
+            ) from exc
+    else:
+        if not hasattr(arr, "data_ptr"):
+            raise RuntimeError(
+                "_cuda_zerocopy: this MLX build neither exports a native kDLCUDA "
+                "DLPack capsule (__dlpack__) nor exposes mx.array.data_ptr(); "
+                "the zero-copy MLX->torch bridge cannot run. Build MLX at the "
+                "kDLCUDA-export commit (6da6a0e4) or a build with data_ptr() "
+                "(MLX #3342)."
+            )
+        capsule = mlx_cuda_array_to_dlpack_capsule(arr)
+        try:
+            t = _torch_from_dlpack(capsule)
+        except Exception as exc:  # noqa: BLE001 - surface where + what failed
+            raise RuntimeError(
+                f"_cuda_zerocopy: torch.utils.dlpack.from_dlpack rejected the "
+                f"MLX-CUDA data_ptr() capsule (device_type={_dlpack_device_type()}, "
+                f"device_id={_cuda_device_id()}, dtype={arr.dtype}, "
+                f"shape={tuple(int(d) for d in arr.shape)}): "
+                f"{type(exc).__name__}: {exc}"
+            ) from exc
+    if not t.is_cuda:
+        raise RuntimeError(
+            "_cuda_zerocopy: torch imported the MLX-CUDA DLPack as a non-CUDA "
+            f"tensor (device={t.device}); the zero-copy bridge requires a CUDA "
+            "device view."
+        )
+    del torch
+    return t
diff --git a/cppmega_mlx/nn/_tilelang/_sparse_mla_v32_fused.py b/cppmega_mlx/nn/_tilelang/_sparse_mla_v32_fused.py
@@ -28,13 +28,22 @@
 
 gb10 / sm_121 note
 ------------------
-The v32 fwd/bwd tiles request >48 KB dynamic shared memory. On gb10/sm_121 the
-TileLang/TVM runtime's ``cuFuncSetAttribute(MAX_DYNAMIC_SHARED_SIZE_BYTES)``
-opt-in is rejected by the driver above ~48 KB even though the device reports a
-99 KB opt-in carve-out, so the fused kernel does not yet run there (it lowers +
-compiles for sm_121a; only the runtime smem opt-in is blocked). On H100 (sm_90)
-/ B200 (sm_100) the carve-out fits and the fused path runs. This is surfaced
-loudly when forced, never silently degraded.
+RESOLVED (tilelang 823c807c): the v32 fwd/bwd tiles now FIT the gb10/sm_121
+99 KiB dynamic-smem carve-out via the re-tiled GB10 variant (``gb10=True``:
+block_I=16, num_stages=1, aggressive shared-memory merge, Hopper TMA-lower
+disabled, ``O_shared`` dropped, mask-in-shared). MEASURED on real sm_121a ptxas
+(CUDA 13.3): fwd 93.0 KiB, bwd <= 99 KiB; fwd cos 0.998, bwd dq/dkv cos 1.000 vs
+the upstream reference. The earlier >48 KB ``cuFuncSetAttribute`` rejection was a
+codegen/argbinder smem-emission bug, fixed in the merge-codegen-reorg branch.
+
+This wrapper forwards ``gb10`` (auto-detect by default, overridable via
+``CPPMEGA_SPARSE_MLA_V32_GB10``) so the kernel emits the GB10-fitting variant on
+sm_12x and the original Hopper kernel on sm_90/sm_100. MLX-CUDA q/kv/indices are
+fed to the torch-backend kernel interfaces zero-copy via DLPack when
+``CPPMEGA_TILELANG_CUDA_ZEROCOPY=1`` (kDLCUDA ``DLManagedTensor`` over
+``mx.array.data_ptr()``, no host roundtrip); outputs come back from torch (MLX
+cannot import a CUDA DLPack, so the torch->MLX writeback is a host bounce — an
+MLX limitation, surfaced honestly, not a silent degrade of the input path).
 """
 
 from __future__ import annotations
@@ -77,21 +86,71 @@ def _torch():
 
 
 def _mlx_to_torch_cuda(arr: mx.array):
-    """Materialize an MLX array as a contiguous CUDA torch tensor (no host copy
-    when a CUDA->CUDA DLPack bridge is available; otherwise via DLPack)."""
+    """Materialize an MLX array as a contiguous CUDA torch tensor.
+
+    When the zero-copy DLPack escape hatch is enabled
+    (``CPPMEGA_TILELANG_CUDA_ZEROCOPY=1``) the MLX-CUDA device buffer is handed to
+    torch via a kDLCUDA ``DLManagedTensor`` capsule with NO host roundtrip — the
+    returned tensor is a real device view of the MLX allocation. Otherwise we use
+    the numpy-host-roundtrip eager bridge (``_cuda_eager._mlx_to_torch_cuda``).
+
+    RULE #1: the zero-copy bridge RAISES with where+what on any failure; we never
+    silently fall back to the eager host-copy bridge while claiming zero-copy.
+    """
 
     torch = _torch()
-    from cppmega_mlx.nn._tilelang._cuda_eager import _mlx_to_torch_cuda as _bridge
+    if _zerocopy_enabled():
+        from cppmega_mlx.nn._tilelang._cuda_zerocopy import (
+            mlx_cuda_array_to_torch_tensor as _zc_bridge,
+        )
 
-    t = _bridge(arr)
+        t = _zc_bridge(arr)
+    else:
+        from cppmega_mlx.nn._tilelang._cuda_eager import _mlx_to_torch_cuda as _bridge
+
+        t = _bridge(arr)
     if not t.is_cuda:
         raise RuntimeError(
             "_sparse_mla_v32_fused: MLX->torch bridge did not yield a CUDA tensor; "
             "the fused v32 Sparse-MLA path requires CUDA buffers."
         )
+    # ``.contiguous()`` is a no-op (returns the same storage) for the row-major
+    # zero-copy view; it only copies if MLX handed us a strided buffer.
     return t.contiguous()
 
 
+def _zerocopy_enabled() -> bool:
+    """Return whether the zero-copy CUDA DLPack escape hatch is env-gated ON."""
+
+    from cppmega_mlx.nn._tilelang._cuda_zerocopy import zerocopy_enabled
+
+    return zerocopy_enabled()
+
+
+def _gb10_flag() -> bool | None:
+    """Resolve the ``gb10`` arch flag forwarded to the v32 fwd/bwd interfaces.
+
+    The upstream kernels accept ``gb10=None`` (auto-detect sm_12x), ``gb10=True``
+    (force the 99 KiB-fitting GB10 variant: block_I=16, num_stages=1, aggressive
+    smem merge, TMA-lower disabled), or ``gb10=False`` (the original Hopper
+    kernel). We default to auto-detect (``None``) so the kernel itself decides by
+    compute capability; ``CPPMEGA_SPARSE_MLA_V32_GB10`` forces the choice
+    (1/true/on -> True, 0/false/off -> False) for A/B testing or non-sm_12x hosts.
+    """
+
+    raw = os.environ.get("CPPMEGA_SPARSE_MLA_V32_GB10", "").strip().lower()
+    if raw == "":
+        return None
+    if raw in {"1", "true", "yes", "on"}:
+        return True
+    if raw in {"0", "false", "no", "off"}:
+        return False
+    raise ValueError(
+        "_sparse_mla_v32_fused: CPPMEGA_SPARSE_MLA_V32_GB10 must be one of "
+        f"1/true/on or 0/false/off (or unset for auto-detect); got {raw!r}."
+    )
+
+
 def _torch_cuda_to_mlx(tensor, dtype: mx.Dtype) -> mx.array:
     from cppmega_mlx.nn._tilelang._cuda_eager import _torch_cuda_to_mlx as _bridge
 
@@ -204,6 +263,7 @@ def v32_fused_apply(
             kv_t.contiguous(),
             idx_t.contiguous(),
             sm_scale=float(sm_scale),
+            gb10=_gb10_flag(),
         )
         torch.cuda.synchronize()
     except Exception as exc:
@@ -256,6 +316,7 @@ def v32_fused_bwd(
             idx_t,
             lse_t,
             sm_scale=float(sm_scale),
+            gb10=_gb10_flag(),
         )
         torch.cuda.synchronize()
     except Exception as exc:
@@ -268,5 +329,7 @@ def v32_fused_bwd(
         ) from exc
 
     dq = _torch_cuda_to_mlx(dq_t, mx.bfloat16)
-    dkv = _torch_cuda_to_mlx(dkv_t.astype(torch.bfloat16), mx.bfloat16)
+    # ``sparse_mla_bwd``'s postprocess emits dkv as a torch fp32 tensor; cast to
+    # bf16 with torch's ``.to`` (NOT MLX/numpy ``.astype``) before the writeback.
+    dkv = _torch_cuda_to_mlx(dkv_t.to(torch.bfloat16), mx.bfloat16)
     return dq, dkv
diff --git a/docs/FUSED-PIPELINE-ROADMAP.md b/docs/FUSED-PIPELINE-ROADMAP.md
@@ -123,15 +123,15 @@ 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): WIRED + ENV-GATED; BLOCKED ON gb10/sm_121 (NOT atomics — driver smem opt-in).**
-- **Wiring landed (default OFF):** new `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`. Gate OFF = existing reference path, unchanged (50 structural tests still green; the one local failure is a pre-existing tvm_ffi circular-import in the Metal FP8 producer, reproduced on unmodified `main`). RULE #1: when gated ON + `force_path_c`, every fused failure RAISES with where+what — no silent reference fallback.
-- **gb10 smoke-test verdict (single idle run, free>90G, cleaned after):** the fused fwd/bwd does **NOT** run on gb10/sm_121 today. The blocker is **NOT** `atomic_addx4` and **NOT** a true smem overflow — it is the TileLang/TVM runtime's dynamic-shared-memory opt-in. Three findings, in order hit:
-  1. **`T.dynamic` symbolic-shape lowering crash** (gb10 tilelang @ upstream `main` `c0a6fe5a`): `LowerDeviceKernelLaunch` raises `Check failed: new_data_expr->IsInstance<VarNode>() ... backing allocation must be a tirx::Var` on the `T.dynamic("batch/seq_len")` buffers — the `tir→tirx` migration regression (roadmap §6 grid/symbolic risk, confirmed). Static shapes dodge it.
-  2. **smem overflow at default tiling:** static fwd compiles for `-arch=sm_121a` then `ptxas error: Entry function 'main_kernel' uses too much shared data (0x38800=231424 bytes, 0x18c00=101376 max)` — the H100/B200-class 228 KB tile vs gb10's 99 KB carve-out.
-  3. **the real wall — runtime smem opt-in rejected:** re-tiled down to ~56 KB, the kernel fails at `cuda_module.cc:218` `cuFuncSetAttribute(CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES, 57344)` → `Failed to set the allowed dynamic shared memory size`. **Bisected with a trivial kernel: the driver accepts ≤50000 B but rejects ≥51200 B**, i.e. the >48 KB opt-in carve-out is **unusable** on this gb10 stack even though the device reports `shared_memory_per_block_optin=101376`. The fused tile's fixed `H_per_block×512` + `BI×512` bf16 buffers cannot fit under ~48 KB (min viable `block_I=32` already needs ~56 KB; `block_I=16` violates a GEMM `warp_col_tiles>8` constraint). So **no valid fused config fits the gb10 runtime ceiling.**
-- **`atomic_addx4` on sm_121: UNDETERMINED** — the bwd never reached the atomic scatter because the fwd (and any dyn-smem kernel) is blocked at the smem opt-in first. NOT verified working, NOT verified broken.
-- **Parity:** the fused fwd+bwd-vs-torch parity is the upstream v32 example's own `assert_tensors_similar` (passes on sm_90). End-to-end fused-vs-reference parity could **not** be measured on gb10 (kernel can't run) — honestly reported, not faked.
-- **Net:** on **H100 (sm_90) / B200 (sm_100)** the 99 KB+ carve-out fits and the gated fused path should run; on **gb10 the fused v32 MLA is blocked** on the runtime dynamic-smem opt-in (a TVM `cuda_module.cc` / driver issue), not the kernel ISA. Fallback per §3.1 rank-5 (NSA `deepseek_nsa/*_bwd.py`) is the next probe — but NSA tiles likely hit the same >48 KB opt-in wall, so the durable gb10 fix is patching the TVM runtime smem-opt-in call (or re-tiling every dyn-smem kernel under 48 KB), tracked as a gb10 follow-on.
+**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).**
+
+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.
+
+- **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.
 
 ### 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.
diff --git a/scripts/verify_sparse_mla_v32_pathc_e2e.py b/scripts/verify_sparse_mla_v32_pathc_e2e.py
