feat(v4-path-d): generalize GDN Triton-frontend Path D to KDA-parity + runnable contract

GDN Path D (FLA Triton -> TileLang -> Metal) was fixed-prefill-only (Bx64xH1xK64/V32), test_only.
Generalized to the full KDA-equivalent surface:
- _compile_gdn_runtime_stages parametrized (total_tokens/num_sequences/num_chunks/h/hv/k/v/varlen/
  scale/initial+final state/topology), mirroring _compile_kda_runtime_stages.
- gdn_fwd_runtime_call reads shapes from inputs (HV%H==0, K<=256), custom scale, init/final state,
  packed varlen; added GDN scalar-gate chunk-local cumsum.
- Backward = reuse Path B's fused recurrent Metal VJP (gdn_apply_path_d) - same (q,k,v,beta,g)
  signature, matches Path A fwd ~1e-8, runs M1+ (lowering FLA chunk_gated_delta_rule_bwd would hit
  the same M5 GEMM wall + need new OP_TABLE coverage). 5 finite grads.
- Fixed a varlen ABI bug (kkt/recompute hardcoded cu_seqlens/chunk_indices len=1 -> TVM-FFI binding
  error); _gdn_varlen_meta_lengths.
- New parity harness scripts/v4_gdn_path_d_parity.py (8 production-like shapes vs Path A FLA-naive).

Contract: v4.gdn path_d test_only/blocked -> runnable_if_available/candidate, matching the already
shipped KDA row (both compile+ABI green, numeric launch M5-gated). Default dispatch UNCHANGED:
without the explicit unsafe-import env + local FLA/triton checkouts, Path D fail-closes to
unavailable -> auto-pick stays path_c.

Verified on M4 Max: all 8 harness shapes backend_available=True (lowering+compile+ABI green);
backward finite. NUMERIC LAUNCH PARITY IS M5-GATED - the FLA chunk kernels lower to TileLang
mpp::tensor_ops::matmul2d (Metal 4 / Apple M5); M4 launch fails 'undeclared mpp' (KDA Path D hits
the identical wall). Capture green numeric parity on M5 via the harness. Tests 65 passed (failing
launch-smoke tests pre-existing on pristine main, same M5 blocker); v3 untouched.

Author: jeweldave

cppmega_v4/_tilelang/linear_attention_path_d_real.py

@@ -533,6 +533,26 @@ def _full_arg(extent: int) -> Tuple[int]:
     return (max(int(extent), 1),)
 
 
+def _gdn_varlen_meta_lengths(
+    constexprs: Dict[str, Any],
+    b: int,
+) -> Tuple[int, int]:
+    """Return ``(cu_seqlens_len, chunk_indices_len)`` honoring IS_VARLEN.
+
+    When packed varlen is active the kkt/recompute/chunk_o kernels receive the
+    real cu_seqlens (length N+1) and chunk_indices (2 ints per chunk); under
+    fixed-length they are unused single-element placeholders.
+    """
+
+    t = _runtime_t(constexprs)
+    bt = int(constexprs["BT"])
+    if bool(constexprs.get("IS_VARLEN", False)):
+        n = _runtime_n(constexprs, b)
+        nt = _runtime_num_chunks(constexprs, _num_chunks(t, bt))
+        return _runtime_cu_len(constexprs, n + 1), nt * 2
+    return 1, 1
+
+
 def _gdn_kkt_arg_buffer_shapes(
     constexprs: Dict[str, Any],
     grid: Optional[Tuple[int, ...]],
@@ -543,13 +563,14 @@ def _gdn_kkt_arg_buffer_shapes(
     k = int(constexprs["K"])
     bt = int(constexprs["BT"])
     b = _batch_from_grid(grid, hv)
+    cu_len, chunk_indices_len = _gdn_varlen_meta_lengths(constexprs, b)
     return {
         0: _full_arg(b * t * h * k),      # k
         1: _full_arg(b * t * hv),         # g
         2: _full_arg(b * t * hv),         # beta
         3: _full_arg(b * t * hv * bt),    # A
-        4: _full_arg(1),                  # cu_seqlens, unused when IS_VARLEN=False
-        5: _full_arg(1),                  # chunk_indices, unused when IS_VARLEN=False
+        4: _full_arg(cu_len),             # cu_seqlens
+        5: _full_arg(chunk_indices_len),  # chunk_indices
     }
 
 
@@ -564,6 +585,7 @@ def _gdn_recompute_arg_buffer_shapes(
     v = int(constexprs["V"])
     bt = int(constexprs["BT"])
     b = _batch_from_grid(grid, hv)
+    cu_len, chunk_indices_len = _gdn_varlen_meta_lengths(constexprs, b)
     return {
         0: _full_arg(b * t * h * k),      # k
         1: _full_arg(b * t * hv * v),     # v
@@ -572,8 +594,8 @@ def _gdn_recompute_arg_buffer_shapes(
         4: _full_arg(b * t * hv * v),     # u
         5: _full_arg(b * t * hv * bt),    # A
         6: _full_arg(b * t * hv),         # g
-        7: _full_arg(1),                  # cu_seqlens
-        8: _full_arg(1),                  # chunk_indices
+        7: _full_arg(cu_len),             # cu_seqlens
+        8: _full_arg(chunk_indices_len),  # chunk_indices
     }