test(v4-path-c): lock in GDN+KDA Path C real-kernel parity (unblocked by ffi-NULL fix)

The ffi-NULL fix (tilelang 168877b0) unblocked GDN Path C with NO prim completion needed -- the
scaffold's recurrence math was already correct (structural twin of the working KDA Path C kernel).
Both GDN + KDA Path C now DISPATCH the real @T.prim_func Metal kernel and MATCH the FLA-naive
reference (rel_err 1.9e-7..4.8e-7, allclose True across production shapes, L2-normalized q/k).

Tests (test-only, no production code changed):
- +test_kda_path_c_real_kernel_parity_with_path_a (the KDA analog the suite lacked).
- Fixed stale test_gdn_dispatch_each_path_runs[path_e]: it fed a sign-unconstrained gate (g>0
  possible), physically invalid for GDN (decay=exp(g)<=1); Path E correctly RAISED per RULE #1.
  Changed to a valid non-positive gate (pre-existing failure on origin, not introduced here).

The 2 masking tests (test_path_c_forced_returns_valid_output GDN+KDA) now pass via REAL Path C
(spy-confirmed, not Path A). Full GDN+KDA suite 109 passed/2 skip. RULE #1 verified: runtime crash
in Path C raises "Refusing to silently fall back to Path A" even with allow_fallback=True.

Author: jeweldave

tests/v4/test_kda_paths.py

@@ -9,6 +9,7 @@
 from cppmega_v4._tilelang.kda_path_b import kda_forward_path_b
 from cppmega_v4._tilelang.kda_path_c import (
     _path_c_runtime_status as kda_path_c_runtime,
+    _tilelang_importable as kda_tilelang_importable,
 )
 from cppmega_v4._tilelang.kda_path_d import (
     _fla_kda_chunk_importable,
@@ -196,6 +197,37 @@ def test_kda_path_c_fallback_matches_path_a():
     np.testing.assert_array_equal(np.array(o_disp), np.array(o_ref))
 
 
+@pytest.mark.skipif(
+    not kda_tilelang_importable()[0],
+    reason="tilelang not importable in this env",
+)
+def test_kda_path_c_real_kernel_parity_with_path_a():
+    """Real @T.prim_func KDA Path C parity vs the FLA-naive reference.
+
+    Uses L2-normalised q/k (the production conditioning for KDA — q and k are
+    always normalised in the real architecture) so the delta-rule recurrence
+    stays bounded and fp32 reassociation between the Metal kernel and the MLX
+    reference is at the 1e-6 level rather than tripping on an exploded state.
+    """
+    from cppmega_v4._tilelang.kda_path_c import _kda_fwd_path_c_call
+
+    B, T, H, HV, K, V = 1, 16, 2, 4, 64, 64
+    rng = np.random.default_rng(73)
+
+    def _l2n(x):
+        return x / (mx.linalg.norm(x, axis=-1, keepdims=True) + 1e-6)
+
+    q = _l2n(mx.array(rng.standard_normal((B, T, H, K)).astype(np.float32)))
+    k = _l2n(mx.array(rng.standard_normal((B, T, H, K)).astype(np.float32)))
+    v = mx.array(rng.standard_normal((B, T, HV, V)).astype(np.float32))
+    g = mx.array(-rng.uniform(0.01, 0.2, (B, T, HV, K)).astype(np.float32))
+    beta = mx.array(rng.uniform(0.1, 0.9, (B, T, HV)).astype(np.float32))
+    o_c, _ = _kda_fwd_path_c_call(q, k, v, g, beta)
+    o_a, _ = naive_recurrent_kda(q, k, v, g, beta)
+    mx.eval(o_c, o_a)
+    np.testing.assert_allclose(np.array(o_c), np.array(o_a), atol=1e-4, rtol=1e-3)
+
+
 # ----- Path D (Triton frontend) -----
 
 

tests/v4/test_path_dispatch.py

@@ -136,12 +136,18 @@ def test_gdn_dispatch_returns_same_as_path_a():
 
 @pytest.mark.parametrize("path", ["path_a", "path_b", "path_c", "path_d", "path_e"])
 def test_gdn_dispatch_each_path_runs(path):
-    """Each forced path must return finite, correctly-shaped output."""
+    """Each forced path must return finite, correctly-shaped output.
+
+    The GDN gate is parameterised as decay = exp(g) <= 1, so g must be <= 0
+    (a g>0 amplifying gate has no real pre-image under softplus_inverse(-g),
+    and Path E correctly RAISES per RULE #1 rather than silently clamping).
+    Use a physically-valid non-positive gate so every forced path runs.
+    """
     q = mx.random.normal((1, 4, 2, 4))
     k = mx.random.normal((1, 4, 2, 4))
     v = mx.random.normal((1, 4, 2, 4))
-    beta = mx.random.normal((1, 4, 2))
-    g = mx.random.normal((1, 4, 2)) * 0.1
+    beta = mx.sigmoid(mx.random.normal((1, 4, 2)))
+    g = -mx.abs(mx.random.normal((1, 4, 2)) * 0.1)
     o, _ = gated_delta_recurrent_dispatch(q, k, v, beta, g, path=path)
     assert o.shape == (1, 4, 2, 4)
     assert not bool(mx.any(mx.isnan(o)).item())