feat(v4-path-d): recompute_w_u threadgroup tiling + correct Metal-4 contract (not M5-gated)

recompute_w_u materialized ~27K elems/thread in thread-private arrays -> pipeline-state
'exceeds available stack space'. Tile through threadgroup memory (BK=32/BV=16 constexprs,
GDN + KDA) -> recompute_w_u now compiles+launches on M4. Contract reword: Path D runs on
Metal 4 (macOS 26+) incl Apple M4 -- the gate is OS/SDK capability, NOT M5 silicon (my
earlier 9c7d14f 'M5-gated' reason was wrong). Note: full Path D parity still blocked by a
separate pre-existing launch-ABI num_args mismatch (make_packed_api), independent of the
now-correct matmul2d.

Author: jeweldave

cppmega_v4/_tilelang/path_d_runtime_adapter.py

@@ -37,6 +37,16 @@
 GDN_RUNTIME_BT = 32
 GDN_CHUNK_H_METAL_SAFE_BV = 16
 GDN_FIXED_CHUNK_O_BV = 16
+# recompute_w_u (shared by GDN and KDA) materializes several BT*BK + BT*BV
+# THREAD-PRIVATE accumulators (b_w, b_u, b_kb, b_vb, b_A, ...). With the FLA
+# default BK=BV=64 these are ~4096 elements/thread (~27K total), which makes
+# newComputePipelineStateWithFunction fail with "Compute function exceeds
+# available stack space" on Apple Metal. Tiling the K/V reduction through
+# smaller blocks shrinks each thread-private tile so the pipeline-state
+# creation fits Metal's per-thread stack budget. This is an occupancy/tiling
+# limit (it would fail on M5 too), independent of the Metal-4 language gate.
+GDN_RECOMPUTE_W_U_METAL_SAFE_BK = 32
+GDN_RECOMPUTE_W_U_METAL_SAFE_BV = 16
 KDA_FIXED_T = 64
 KDA_FIXED_H = 1
 KDA_FIXED_HV = 1
@@ -1472,6 +1482,11 @@ def compile_kda_recompute_w_u_artifact(
             "K": int(k_dim),
             "V": int(v_dim),
             "BT": int(bt),
+            # Tile K/V through Metal-safe blocks so the per-thread WY-recompute
+            # accumulators fit the pipeline-state stack budget (shared root
+            # cause with GDN recompute_w_u -> "exceeds available stack space").
+            "BK": min(int(k_dim), GDN_RECOMPUTE_W_U_METAL_SAFE_BK),
+            "BV": min(int(v_dim), GDN_RECOMPUTE_W_U_METAL_SAFE_BV),
             "IS_VARLEN": bool(is_varlen),
         }
     )
@@ -1849,7 +1864,15 @@ def _compile_gdn_runtime_stages(
         "IS_VARLEN": bool(is_varlen),
     }
     kkt_constexprs = {**runtime_private, **shared_dims}
-    recompute_constexprs = {**runtime_private, **shared_dims, "V": int(v_dim)}
+    recompute_constexprs = {
+        **runtime_private,
+        **shared_dims,
+        "V": int(v_dim),
+        # Tile the K/V reduction through Metal-safe blocks so the per-thread
+        # recompute_w_u accumulators fit the pipeline-state stack budget.
+        "BK": min(int(k_dim), GDN_RECOMPUTE_W_U_METAL_SAFE_BK),
+        "BV": min(int(v_dim), GDN_RECOMPUTE_W_U_METAL_SAFE_BV),
+    }
     chunk_h_constexprs = {
         **runtime_private,
         "H": int(h_heads),

scripts/path_sanity_guard.py

@@ -184,9 +184,13 @@ class PathContract:
         "initial/final recurrent state, packed varlen) with a fused Path B "
         "backward, mirroring the KDA Path D row. See "
         "reports/raw/v4_gdn_path_d_parity.json; numeric-launch parity vs the "
-        "Path A FLA-naive reference is captured on Apple M5 hardware (the "
-        "TileLang mpp::tensor_ops::matmul2d cooperative-tensor GEMM the FLA "
-        "kernels lower to requires Metal 4 / M5), identical to KDA Path D.",
+        "Path A FLA-naive reference. The TileLang mpp::tensor_ops::matmul2d "
+        "cooperative-tensor GEMM the FLA kernels lower to runs on Metal 4 "
+        "(macOS 26+) and is verified on Apple M4 (NOT M5-gated): the gate is "
+        "the OS/SDK Metal-4 capability, not the silicon generation. "
+        "recompute_w_u is tiled through Metal-safe BK/BV blocks so its "
+        "per-thread accumulators fit the pipeline-state stack budget, "
+        "identical to KDA Path D.",
     ),
     PathContract(
         "v4.gdn",