fix(v4-path-d): zero-init kkt_solve output buffer -> Path D NaN gone (0 NaN, 0 deadlock)

The Path D NaN (cells with chunks/seq>=3 or multi-head, only when mx.eval(FLA-naive ref) runs before
Path D) was NOT the sync hazard the handoff doc claimed -- every device-event/commit attempt
deadlocked because they solved a non-problem. Real root cause: kkt_solve writes only the
lower-triangular (T_chunk x T_chunk) region of its per-chunk output `a`; the strictly-upper lanes keep
whatever mx.empty allocates. Right after the ref eval MLX frees the reference's NaN-containing fp32
intermediates and mx.empty RECYCLES that freed NaN memory into `a`; downstream chunk GEMMs read the
full matrix (0*NaN=NaN) -> pipeline poison. Without the ref-eval the recycled memory is clean -> the
timing dependence.

Fix: mx.empty -> mx.zeros for `a` (one line). No device event / no command-buffer commit -> no
deadlock. tilelang-side changes reverted (fix entirely in cppmega).

M4 sequential harness (mx.eval(ref) before Path D): 8/8 launched, 5/8 parity-ok, 0 NaN, 0 deadlock,
completes cleanly. 3 originally-NaN cells now finite (varlen_64_32 fully parity-green); remaining
t128/mh mismatch is deterministic fp16 chunk-GEMM precision (separate, not the NaN). No Path C
regression (mamba3 path_c 89, sparse_mla_fp8+#6 65, metal-sync 12). RULE #1: single correct path.

Author: jeweldave

cppmega_v4/_tilelang/path_d_runtime_adapter.py

@@ -2416,7 +2416,18 @@ def gdn_fwd_runtime_call(
         )
     )
 
-    a = mx.empty((b * total_tokens * hv_heads * GDN_RUNTIME_BT,), dtype=mx.float16)
+    # kkt_solve writes only the lower-triangular (T_chunk x T_chunk) region of
+    # its per-chunk output; the strictly-upper lanes are structurally unused by
+    # recompute_w_u but MUST be finite, because downstream chunk GEMMs multiply
+    # the full matrix (0 * NaN == NaN). mx.empty leaves those lanes as whatever
+    # freed allocation MLX recycles -- and right after a heavy fp32 reference
+    # (the FLA-naive parity ref) that recycled memory contains NaN, which then
+    # poisons the whole pipeline (observed: kkt output a == NaN -> y == NaN, only
+    # when mx.eval(ref) runs first and reshuffles allocations). Zero-initialize a
+    # so the unwritten lanes are a hard 0, never recycled NaN. The fully-written
+    # buffers below stay mx.empty. RULE: never feed a partially-written buffer's
+    # uninitialized lanes into a kernel that reads the whole tensor.
+    a = mx.zeros((b * total_tokens * hv_heads * GDN_RUNTIME_BT,), dtype=mx.float16)
     w = mx.empty((b * total_tokens * hv_heads * k_dim,), dtype=mx.float16)
     u = mx.empty((b * total_tokens * hv_heads * v_dim,), dtype=mx.float16)
     v_new = mx.empty((b * total_tokens * hv_heads * v_dim,), dtype=mx.float16)