feat(v4-path-e): harden vendored mlx-lm Path E - fail-close guards, eligibility, provenance, bench

Path E (vendored mlx-lm GDN/KDA gated_delta + FP8 dequant) was runnable but silently traded
correctness for speed. Hardened:
- Amplifying-gate fail-close: the GDN adapter silently clamped (g_clamped=min(g,-1e-6)) on
  amplifying gates (g>0/decay>1, which upstream cannot represent) -> now raises PathEUnavailable
  so the dispatcher falls back to B/A. New _path_e_eligibility.py = single source of truth
  (gate-sign + shape).
- Eligibility in status: GDN/KDA Path E status reflects Dk%32==0 & Dv%4==0 (+ gate sign) so
  auto_pick SKIPS E instead of dropping into the slow upstream-ops path. KDA adapter silent
  slow-path fallback removed (fails closed; KDA has no gate-sign constraint).
- Relabel: lightning_indexer_fp8 no longer self-labels "path E" (pure-MLX + vendored dequant).
- Provenance: VENDORED_MANIFEST.json pins PR #1217/#1224 + per-file sha256; check_vendored_drift
  wired into path_sanity_guard (vendored_snapshot_drift); test + reports/VENDORED_PROVENANCE.md.

Metal verified: eligible->E, g>0->fail-close to A (output==A exactly), ineligible shape->skip,
drift detected. Bench: fused Metal E-VJP 4.1-5.7x faster than Path B fused bwd -> recommend
E-VJP as GDN bwd default when eligible. Tests 83 passed/2 skip on surface; tests/v4 3053 passed
(18 pre-existing fails); v3 untouched.

Author: jeweldave

cppmega_v4/_tilelang/kda_paths.py

@@ -204,13 +204,42 @@ def _path_e_status() -> PathStatus:
         path="path_e", available=True,
         reason=(
             "vendored mlx-lm gated_delta vectorised-gate Metal kernel "
-            "(Dk%32==0 & Dv%4==0; smaller dims fall back to the pure-ops "
-            "reference). The same kernel as GDN Path E — KDA hits the "
-            "g.ndim==4 branch in gated_delta_kernel."
+            "(fast kernel for Dk%32==0 & Dv%4==0; fails closed for smaller "
+            "dims so the dispatcher falls back to Path B/A instead of the slow "
+            "pure-ops reference). The same kernel as GDN Path E — KDA hits the "
+            "g.ndim==4 branch in gated_delta_kernel. No gate-sign constraint "
+            "(KDA passes g_decay=exp(g) directly)."
         ),
     )
 
 
+def _path_e_status_for_inputs(*args, **kwargs) -> PathStatus:
+    """Input-aware KDA Path E status: importability + shape eligibility.
+
+    KDA has no gate-sign constraint, only the fast-kernel shape (Dk%32==0 &
+    Dv%4==0). ``auto_pick`` consumes this so it SKIPS Path E for ineligible
+    shapes rather than selecting the slow upstream ops fallback. Best-effort:
+    unparseable inputs fall back to the static status.
+    """
+    base = _path_e_status()
+    if not base.available:
+        return base
+    # KDA signature: (q, k, v, g, beta, ...). q.shape[-1]=Dk, v.shape[-1]=Dv.
+    try:
+        from cppmega_v4.nn._external._path_e_eligibility import kda_eligibility
+
+        q = kwargs.get("q", args[0] if len(args) > 0 else None)
+        v = kwargs.get("v", args[2] if len(args) > 2 else None)
+        if q is None or v is None:
+            return base
+        elig = kda_eligibility(int(q.shape[-1]), int(v.shape[-1]))
+    except Exception:
+        return base
+    if elig.eligible:
+        return base
+    return PathStatus(path="path_e", available=False, reason=elig.reason)
+
+
 def _path_e_call(*args, allow_fallback: bool = True, **kwargs):
     if not _path_e_status().available:
         return _fallback_or_raise(
@@ -272,7 +301,16 @@ def kda_recurrent_dispatch(
     status_overrides: Mapping[PathName, PathStatus] | None = None,
     **kwargs,
 ):
-    """Call the auto-selected KDA backend, falling back to Path A."""
+    """Call the auto-selected KDA backend, falling back to Path A.
+
+    In auto-mode Path E availability is recomputed from the actual shape so
+    ``auto_pick`` SKIPS Path E for ineligible shapes (which would otherwise
+    drop to the slow upstream ops fallback).
+    """
+    effective_overrides = dict(status_overrides) if status_overrides else {}
+    if path is None and "path_e" not in effective_overrides:
+        effective_overrides["path_e"] = _path_e_status_for_inputs(*args, **kwargs)
+    status_overrides = effective_overrides or None
     statuses = kda_path_statuses(status_overrides=status_overrides)
     if path is None:
         path = kda_auto_mode_for_inputs(status_overrides=status_overrides)

cppmega_v4/_tilelang/linear_attention_paths.py

@@ -219,8 +219,9 @@ def _path_e_status() -> PathStatus:
         return PathStatus(
             path="path_e", available=True,
             reason=(
-                "vendored mlx-lm PR #1217 gated_delta_update (Metal kernel "
-                "for Dk%32==0 & Dv%4==0; ops fallback otherwise)"
+                "vendored mlx-lm PR #1217 gated_delta_update (fast Metal kernel "
+                "for Dk%32==0 & Dv%4==0 and gate g<=0; fails closed otherwise "
+                "so the dispatcher falls back to Path B/A)"
             ),
         )
     except Exception:
@@ -235,6 +236,38 @@ def _path_e_status() -> PathStatus:
         )
 
 
+def _path_e_status_for_inputs(*args, **kwargs) -> PathStatus:
+    """Input-aware Path E status: combine importability with eligibility.
+
+    Path E is only *truly* available for a concrete call when (a) the adapter
+    imports AND (b) the gate is representable (g<=0; no amplifying gate) AND
+    (c) the shape hits the fast Metal kernel. ``auto_pick`` consumes this so it
+    SKIPS Path E (falls through to D/A) rather than selecting a path that would
+    silently clamp the gate or drop to the slow upstream ops fallback.
+
+    Best-effort: if inputs cannot be parsed (e.g. unusual call shape) we fall
+    back to the static status so behaviour is never worse than before.
+    """
+    base = _path_e_status()
+    if not base.available:
+        return base
+    # Expected signature: (q, k, v, beta, g, ...). Probe g + dims defensively.
+    try:
+        from cppmega_v4.nn._external._path_e_eligibility import gdn_eligibility
+
+        q = kwargs.get("q", args[0] if len(args) > 0 else None)
+        v = kwargs.get("v", args[2] if len(args) > 2 else None)
+        g = kwargs.get("g", args[4] if len(args) > 4 else None)
+        if q is None or v is None or g is None:
+            return base
+        elig = gdn_eligibility(g, int(q.shape[-1]), int(v.shape[-1]))
+    except Exception:
+        return base
+    if elig.eligible:
+        return base
+    return PathStatus(path="path_e", available=False, reason=elig.reason)
+
+
 def _path_e_call(*args, allow_fallback: bool = True, **kwargs):
     status = _path_e_status()
     if not status.available:
@@ -299,7 +332,16 @@ def gated_delta_recurrent_dispatch(
     """Call the auto-selected GDN backend, falling back to Path A.
 
     Same callable signature as ``naive_recurrent_gated_delta_rule``.
+
+    In auto-mode (``path is None``) Path E availability is recomputed from the
+    *actual* inputs (gate sign + shape) so ``auto_pick`` SKIPS Path E for
+    amplifying gates or ineligible shapes — it never selects a path that would
+    silently clamp the gate or drop to the slow upstream ops fallback.
     """
+    effective_overrides = dict(status_overrides) if status_overrides else {}
+    if path is None and "path_e" not in effective_overrides:
+        effective_overrides["path_e"] = _path_e_status_for_inputs(*args, **kwargs)
+    status_overrides = effective_overrides or None
     statuses = linear_attention_path_statuses(status_overrides=status_overrides)
     if path is None:
         path = linear_attention_auto_mode_for_inputs(status_overrides=status_overrides)

cppmega_v4/nn/_external/VENDORED_MANIFEST.json

@@ -0,0 +1,88 @@
+{
+  "_comment": [
+    "Provenance pin for the mlx-lm upstream code vendored into cppmega_v4 Path E",
+    "and the V4 mHC / MTP stack. Each entry records the upstream PR, the path it",
+    "was copied from, whether it is a verbatim snapshot or a cppmega-authored",
+    "derivative, and a sha256 of the vendored file as committed. The drift check",
+    "(_vendored_provenance.py, exercised by tests/v4/test_vendored_provenance.py)",
+    "recomputes these hashes and fails if a vendored file diverges from its",
+    "recorded snapshot, so silent edits to a 'verbatim' copy cannot slip in.",
+    "Hashes are over the file bytes exactly as stored in the repo."
+  ],
+  "schema_version": 1,
+  "upstream_repo": "https://github.com/ml-explore/mlx-lm",
+  "upstream_license": "MIT (Copyright Apple Inc.)",
+  "recorded_at": "2026-05-30",
+  "entries": [
+    {
+      "file": "_mlx_lm_gated_delta_vendored.py",
+      "pr": 1217,
+      "pr_url": "https://github.com/ml-explore/mlx-lm/pull/1217",
+      "upstream_path": "mlx_lm/models/gated_delta.py",
+      "kind": "verbatim",
+      "used_by": "GDN/KDA Path E (gated_delta_update, gated_delta_kernel, gated_delta_ops, compute_g)",
+      "sha256": "18b3a77250b246d46e63d8fa9e8208936ea09bd46b97b036c2ec34edfc0a176b",
+      "notes": "Snapshot includes the PR #1066 Kahan-compensated kv_mem accumulation + 4-way time-loop unroll; this is NOT identical to an older mlx-lm checkout that predates PR #1066."
+    },
+    {
+      "file": "_mlx_lm_fp8_dequant_vendored.py",
+      "pr": 1224,
+      "pr_url": "https://github.com/ml-explore/mlx-lm/pull/1224",
+      "upstream_path": "mlx_lm/models/qwen3_5_moe.py (Model.sanitize FP8 dequant block, generalized)",
+      "kind": "derived",
+      "used_by": "Lightning Indexer FP8 (pure-MLX, NOT Path E) + MoE FP8 loaders",
+      "sha256": "3a48eec1797aee1a9c6d16e819064a045624a3eb62faf69a5670faf1b978f2a8",
+      "notes": "Inline upstream snippet pulled into a standalone dequant_block_fp8 utility; functionally faithful but reorganised, hence 'derived'."
+    },
+    {
+      "file": "_mlx_lm_gated_delta_vjp_metal_vendored.py",
+      "pr": 1217,
+      "pr_url": "https://github.com/ml-explore/mlx-lm/pull/1217",
+      "upstream_path": "mlx_lm/models/gated_delta.py (kernel reused; VJP is cppmega-authored)",
+      "kind": "derived",
+      "used_by": "GDN Path E training backward (fused Metal VJP)",
+      "sha256": "dfe959d0ed50d60d75d93532a80b0b6f769aad74ce39bc00a0770128c9903f4e",
+      "notes": "cppmega-authored fused Metal VJP built on the vendored PR #1217 gated_delta_kernel; upstream has no registered VJP."
+    },
+    {
+      "file": "_mlx_lm_gated_delta_vjp_vendored.py",
+      "pr": 1217,
+      "pr_url": "https://github.com/ml-explore/mlx-lm/pull/1217",
+      "upstream_path": "mlx_lm/models/gated_delta.py (Python-ops reference; VJP is cppmega-authored)",
+      "kind": "derived",
+      "used_by": "GDN Path E training backward (gradient-checkpointed Python reference)",
+      "sha256": "1a5040eb617070f6fbc0712af88cd35104d83a6917a0c5fdfb518dc3a1436b10",
+      "notes": "cppmega-authored chunked Python VJP wrapping the vendored PR #1217 ops; bounds T>=2048 memory."
+    },
+    {
+      "file": "_mlx_lm_hyper_connection_vendored.py",
+      "pr": 1189,
+      "pr_url": "https://github.com/ml-explore/mlx-lm/pull/1189",
+      "upstream_path": "mlx_lm/models/hyper_connection.py",
+      "kind": "verbatim",
+      "used_by": "V4 mHC residual (HyperConnection, HyperHead)",
+      "sha256": "6febc66a50e6504fad2f60b2af51e3eb01a9fb83713b4189d39cf9c499b7b9c7",
+      "notes": "Recorded for completeness; not part of Path E."
+    },
+    {
+      "file": "_mlx_lm_sinkhorn_vendored.py",
+      "pr": 1189,
+      "pr_url": "https://github.com/ml-explore/mlx-lm/pull/1189",
+      "upstream_path": "mlx_lm/models/sinkhorn.py",
+      "kind": "verbatim",
+      "used_by": "V4 mHC Sinkhorn projection (hc_split_sinkhorn)",
+      "sha256": "af2dee5ef7db71c9bc254bae8000ade7e92fd8e5a8367f3bd3bf4b713b1a587f",
+      "notes": "Recorded for completeness; not part of Path E."
+    },
+    {
+      "file": "_mlx_lm_mtp_module_vendored.py",
+      "pr": 990,
+      "pr_url": "https://github.com/ml-explore/mlx-lm/pull/990",
+      "upstream_path": "mlx_lm/models/qwen3_5.py (MTPModule excerpt)",
+      "kind": "derived",
+      "used_by": "V4 MTP speculative-decode module",
+      "sha256": "3131e547b15a43b2538ccfb1c14a275f91f51739008dca4770a28c7fb414444c",
+      "notes": "Excerpt only (MTPModule + decoder layer); the PR's generate/server drivers are out of scope."
+    }
+  ]
+}

cppmega_v4/nn/_external/_path_e_eligibility.py

@@ -0,0 +1,145 @@
+"""Path E eligibility checks — single source of truth.
+
+Path E wraps the vendored mlx-lm ``gated_delta`` Metal kernel (PR #1217). It
+has two HARD limits that, if ignored, silently trade correctness or speed:
+
+1. GATE SIGN (correctness).
+   Upstream parameterises the per-step decay as
+   ``g_decay = exp(-exp(A_log) * softplus(a + dt_bias))`` which is bounded by
+   ``1`` for all real inputs. The GDN adapter recovers our FLA log-decay ``g``
+   via ``a = softplus_inverse(-g)``, which only exists for ``g <= 0`` (decay
+   ``<= 1``). A model whose gate ``g > 0`` (an *amplifying* gate, decay > 1)
+   CANNOT be represented; the historic adapter silently clamped ``g`` to ``0``
+   and lost information. We instead FAIL-CLOSE: Path E reports itself
+   unavailable / raises so the dispatcher falls back to Path B/A, which can
+   represent amplifying gates exactly.
+
+   KDA passes ``g_decay = exp(g)`` straight into the kernel, so KDA itself has
+   no upper bound on the gate — but the kernel multiplies state by a *positive*
+   decay each step, and a NEGATIVE ``g_decay`` (which never arises from
+   ``exp``) would be ill-defined. KDA's amplifying gates (``g > 0``) are
+   therefore representable; only the GDN parameterisation is bounded.
+
+2. SHAPE (silent slow path).
+   The fast Metal kernel only runs when ``Dk % 32 == 0`` and ``Dv % 4 == 0``.
+   For other dims the upstream falls back to a pure-MLX ops reference, which is
+   correct but SLOW. ``auto_pick`` must not select Path E for such shapes —
+   otherwise it picks a "kernel" that is really the slow reference. We record
+   the eligibility in the Path E status so ``auto_pick`` skips it.
+
+These helpers are imported by the GDN/KDA adapters (to fail-close at call
+time) and by the dispatch ``_path_e_status`` functions (to skip E in
+auto-mode). Keeping them here means the rule is defined once.
+"""
+
+from __future__ import annotations
+
+from dataclasses import dataclass
+
+import mlx.core as mx
+
+# Fast Metal kernel constraints (mlx-lm gated_delta.py ``_make_gated_delta_kernel``).
+KERNEL_DK_MULTIPLE = 32
+KERNEL_DV_MULTIPLE = 4
+
+
+@dataclass(frozen=True)
+class PathEEligibility:
+    """Result of a Path E eligibility probe for a concrete call.
+
+    ``eligible`` is True only when Path E can run on the fast Metal kernel
+    AND represents the requested gate exactly. ``reason`` explains a False.
+    """
+
+    eligible: bool
+    reason: str
+    fast_kernel: bool  # shape qualifies for the fast Metal kernel
+    gate_representable: bool  # gate sign is representable by this path
+
+
+def shape_uses_fast_kernel(dk: int, dv: int) -> bool:
+    """True iff (Dk, Dv) hit the fast vendored Metal kernel (not slow ops)."""
+    return (dk % KERNEL_DK_MULTIPLE == 0) and (dv % KERNEL_DV_MULTIPLE == 0)
+
+
+def _gate_has_amplifying(g: mx.array) -> bool:
+    """True iff any element of the log-decay ``g`` is > 0 (decay > 1)."""
+    # tolerance: treat tiny positive jitter as non-amplifying so that
+    # round-trip noise on a g==0 gate does not spuriously fail-close.
+    return bool(mx.any(g.astype(mx.float32) > 1e-6).item())
+
+
+def gdn_eligibility(g: mx.array, dk: int, dv: int) -> PathEEligibility:
+    """Eligibility for the GDN Path E adapter (bounded, decay <= 1).
+
+    Fails closed when the gate is amplifying (``g > 0``) — the upstream
+    parameterisation cannot represent it — or when the shape forces the slow
+    ops fallback.
+    """
+    fast = shape_uses_fast_kernel(dk, dv)
+    amplifying = _gate_has_amplifying(g)
+    gate_ok = not amplifying
+    if amplifying:
+        reason = (
+            "GDN Path E cannot represent an amplifying gate: g>0 (decay>1) is "
+            "outside upstream g_decay=exp(-exp(A_log)*softplus(...)) which is "
+            "bounded by 1; falling back so correctness is not silently clamped"
+        )
+    elif not fast:
+        reason = (
+            f"GDN Path E shape ineligible: Dk={dk} (need %{KERNEL_DK_MULTIPLE}==0) "
+            f"or Dv={dv} (need %{KERNEL_DV_MULTIPLE}==0) would force the slow "
+            "upstream ops fallback, not the fast Metal kernel"
+        )
+    else:
+        reason = "GDN Path E eligible: gate g<=0 and fast Metal kernel shape"
+    return PathEEligibility(
+        eligible=gate_ok and fast,
+        reason=reason,
+        fast_kernel=fast,
+        gate_representable=gate_ok,
+    )
+
+
+def kda_eligibility(dk: int, dv: int) -> PathEEligibility:
+    """Eligibility for the KDA Path E adapter.
+
+    KDA passes ``g_decay = exp(g)`` directly, so any finite gate is
+    representable (decay is always positive). The only hard constraint is the
+    fast-kernel shape; ineligible shapes force the slow ops fallback.
+    """
+    fast = shape_uses_fast_kernel(dk, dv)
+    if not fast:
+        reason = (
+            f"KDA Path E shape ineligible: Dk={dk} (need %{KERNEL_DK_MULTIPLE}==0) "
+            f"or Dv={dv} (need %{KERNEL_DV_MULTIPLE}==0) would force the slow "
+            "upstream ops fallback, not the fast Metal kernel"
+        )
+    else:
+        reason = "KDA Path E eligible: fast Metal kernel shape (gate always representable)"
+    return PathEEligibility(
+        eligible=fast,
+        reason=reason,
+        fast_kernel=fast,
+        gate_representable=True,
+    )
+
+
+class PathEUnavailable(RuntimeError):
+    """Raised by a Path E adapter when it cannot run correctly/fast.
+
+    The dispatcher catches this and falls back to Path B/A. Subclassing
+    ``RuntimeError`` keeps it compatible with the existing
+    ``_dispatch_failure_or_fallback`` handlers.
+    """
+
+
+__all__ = [
+    "KERNEL_DK_MULTIPLE",
+    "KERNEL_DV_MULTIPLE",
+    "PathEEligibility",
+    "PathEUnavailable",
+    "gdn_eligibility",
+    "kda_eligibility",
+    "shape_uses_fast_kernel",
+]