fix(v4-dispatch): classify PathEUnavailable as selection signal, not runtime crash (RULE #1 FIRE)

_path_e_call wrapped op.gated_delta_update in a bare `except Exception` that routed ANY exception
-- including the INTENTIONAL PathEUnavailable (raised when the GDN gate is amplifying g>0, an
availability/selection signal) -- through _dispatch_failure_or_fallback (the always-RAISE crash
path). So forcing path_e with an amplifying gate misclassified "not eligible for this input" as a
kernel crash and tripped the RULE #1 guard.

Fix: catch PathEUnavailable BEFORE the bare except and route it through _fallback_or_raise (mirrors
the static-unavailability routing); keep the bare except for genuine kernel crashes. Same pre-catch
applied to kda_paths.py _path_e_call (KDA adapter raises PathEUnavailable for ineligible shapes).

NOT a RULE #1 weakening: PathEUnavailable is an explicit selection signal (like static
unavailability); _fallback_or_raise still RAISES when allow_fallback=False, and genuine crashes
still route through the always-raise handler. Metal: forced path_e+amplifying+allow_fallback=True
-> falls back to A (no raise); +allow_fallback=False -> raises; genuine crash -> still raises.
test_path_dispatch.py 25 passed.

Author: jeweldave

cppmega_v4/_tilelang/kda_paths.py

@@ -255,8 +255,17 @@ def _path_e_call(*args, allow_fallback: bool = True, **kwargs):
             "path_e", _path_e_status().reason, allow_fallback, *args, **kwargs,
         )
     try:
+        from cppmega_v4.nn._external._path_e_eligibility import PathEUnavailable
         from cppmega_v4.nn._external.mlx_lm_kda_update import kda_update
         return kda_update(*args, **kwargs)
+    except PathEUnavailable as exc:
+        # AVAILABILITY/SELECTION signal (e.g. an ineligible shape), NOT a
+        # runtime crash. Route it like static unavailability above so the
+        # dispatcher falls back to Path B/A (or RAISES when allow_fallback is
+        # False). A genuine kernel crash still hits the always-raise handler.
+        return _fallback_or_raise(
+            "path_e", str(exc), allow_fallback, *args, **kwargs,
+        )
     except Exception as exc:
         return _dispatch_failure_or_fallback(
             "path_e", exc, allow_fallback, *args, **kwargs,

cppmega_v4/_tilelang/linear_attention_paths.py

@@ -280,6 +280,8 @@ def _path_e_status_for_inputs(*args, **kwargs) -> PathStatus:
 
 
 def _path_e_call(*args, allow_fallback: bool = True, **kwargs):
+    from cppmega_v4.nn._external._path_e_eligibility import PathEUnavailable
+
     status = _path_e_status()
     if not status.available:
         return _fallback_or_raise(
@@ -290,6 +292,17 @@ def _path_e_call(*args, allow_fallback: bool = True, **kwargs):
             "cppmega_v4.nn._external.mlx_lm_gated_delta_update"
         )
         return op.gated_delta_update(*args, **kwargs)
+    except PathEUnavailable as exc:
+        # AVAILABILITY/SELECTION signal, NOT a runtime crash. Path E raises
+        # PathEUnavailable (mlx_lm_gated_delta_update.py) when the gate is
+        # amplifying (g>0) or the shape is ineligible — it explicitly cannot
+        # represent this input. That is semantically identical to the static
+        # unavailability routed above, so route it the same way (fall back to
+        # Path B/A, or RAISE when allow_fallback=False). RULE #1 stays intact:
+        # a genuine kernel crash falls through to the always-raise handler.
+        return _fallback_or_raise(
+            "path_e", str(exc), allow_fallback, *args, **kwargs,
+        )
     except Exception as exc:
         return _dispatch_failure_or_fallback(
             "path_e", exc, allow_fallback, *args, **kwargs,