refactor: RULE #1 - kill B1-B3 staged silent fallbacks (22 raises with where+what)

Replace forbidden silent fallbacks with immediate raises; keep genuine feature-absent detection.
- B1 _cuda_eager.py (6 killed, 6 kept): killed the kernel-launch/writeback except->return None
  swallows (sparse_mla/mamba3 fwd+bwd/fp8-quant/fp8-apply/m2rnn cuda-eager) that hid real CUDA
  bugs or silently degraded to MLX-autograd ref; kept the torch-import + cuda_eager_available +
  shape/dtype applicability probes.
- B2 m2rnn_path_c.py (10 killed) + sparse_mla_fp8_path_c.py (5 killed/narrowed): killed the
  except->return None wrapping compile/lowering of the clear path (compile crash != not-applicable;
  availability is probed separately and preserves where+what); narrowed apply_simplify catches to
  ImportError so a real crash in the pass propagates; kept typed availability/Status probes.
- B3 _engine_dispatch.py (1 killed): killed the engine-succeeded-but-MSL-extraction-None ->
  silent legacy-shim substitution (the MSL-instead-of-tvm-ffi fallback) -> now raises, directing to
  explicit CPPMEGA_MLX_TILELANG_ENGINE=shim; kept the genuine engine-absent ImportError branch.

Metal: clear paths still work (m2rnn/mamba3 117 passed/5skip; sparse_mla_fp8 93 passed). Rewrote
test_dispatch..._falls_back_on_extraction_failure -> ..._raises_on_extraction_failure (it was green
ONLY due to the B3 silent shim). Remaining failures pre-existing on origin (DLPack contiguity,
path_b kernel-name). Flagged follow-up: force_path_c gated-off-by-default sites (~2690/3295/3303).

Author: jeweldave

cppmega_mlx/nn/_tilelang/_cuda_eager.py

@@ -264,8 +264,15 @@ def sparse_mla_fwd_cuda_eager(
         )
         kernel(q_t, kv_t, idx_t, sc_t, out_t, lse_t)
         torch.cuda.synchronize()
-    except Exception:
-        return None
+    except Exception as exc:
+        # RULE #1: a CUDA kernel launch/writeback failure here is a real bug in
+        # the selected Sparse-MLA fwd kernel, not a "feature absent" signal. Do
+        # NOT return None to a silent fallback — raise with where+what.
+        raise RuntimeError(
+            f"_cuda_eager.sparse_mla_fwd_cuda_eager: TileLang-CUDA Sparse-MLA "
+            f"fwd kernel launch/writeback failed "
+            f"({type(exc).__name__}: {exc})."
+        ) from exc
 
     out = _torch_cuda_to_mlx(
         out_t.reshape(shapes.batch, shapes.seq_len, shapes.heads, shapes.d_v),
@@ -431,8 +438,15 @@ def mamba3_mimo_fwd_cuda_eager(
         )
         kernel(x_t, B_t, C_t, z_t, A_t, dt_t, D_t, h0_t, y_t, hl_t)
         torch.cuda.synchronize()
-    except Exception:
-        return None
+    except Exception as exc:
+        # RULE #1: a kernel launch/writeback failure is a real bug in the
+        # vendored Mamba3 MIMO fwd kernel, not "feature absent". Raise loud
+        # instead of returning None into a silent pure-MLX-scan fallback.
+        raise RuntimeError(
+            f"_cuda_eager.mamba3_mimo_fwd_cuda_eager: TileLang-CUDA Mamba3 MIMO "
+            f"fwd kernel launch/writeback failed "
+            f"({type(exc).__name__}: {exc})."
+        ) from exc
 
     y = _torch_cuda_to_mlx(y_t, x.dtype)
     h_last = _torch_cuda_to_mlx(hl_t, h0.dtype)
@@ -719,8 +733,16 @@ def mamba3_mimo_bwd_cuda_eager(
             dh0_t,
         )
         torch.cuda.synchronize()
-    except Exception:
-        return None
+    except Exception as exc:
+        # RULE #1: a kernel launch/writeback failure is a real bug in the
+        # vendored Mamba3 MIMO bwd kernel, not "feature absent". Raise loud
+        # instead of returning None into a silent MLX-autograd reference VJP
+        # fallback (which would mask a broken gradient kernel).
+        raise RuntimeError(
+            f"_cuda_eager.mamba3_mimo_bwd_cuda_eager: TileLang-CUDA Mamba3 MIMO "
+            f"bwd kernel launch/writeback failed "
+            f"({type(exc).__name__}: {exc})."
+        ) from exc
 
     # Reduce the per-(b,h,p) lane partials over the P axis (matches the Metal
     # bf16-snapshot reduction axes exactly).
@@ -864,8 +886,16 @@ def fp8_per_token_quant_cuda_eager(x: mx.array) -> tuple[mx.array, mx.array] | N
         scale_t = torch.zeros(rows, dtype=torch.float32, device="cuda")
         kernel(x_t, fp8_t, scale_t)
         torch.cuda.synchronize()
-    except Exception:
-        return None
+    except Exception as exc:
+        # RULE #1: a kernel launch/writeback failure is a real bug in the FP8
+        # per-token quant producer kernel, not "feature absent". Raise loud
+        # instead of returning None (the caller would otherwise re-raise a
+        # generic guard that hides the real producer-kernel error).
+        raise RuntimeError(
+            f"_cuda_eager.fp8_per_token_quant_cuda_eager: TileLang-CUDA per-token "
+            f"FP8 quant kernel launch/writeback failed "
+            f"({type(exc).__name__}: {exc})."
+        ) from exc
 
     fp8 = _torch_cuda_to_mlx(fp8_t.reshape(x.shape), mx.uint8)
     scale = _torch_cuda_to_mlx(scale_t.reshape(x.shape[:-1]), mx.float32)
@@ -1017,8 +1047,16 @@ def fp8_sparse_mla_apply_cuda_eager(
             lse_t,
         )
         torch.cuda.synchronize()
-    except Exception:
-        return None
+    except Exception as exc:
+        # RULE #1: a kernel launch/writeback failure is a real bug in the
+        # prepared-buffer FP8 Sparse-MLA apply kernel, not "feature absent".
+        # Raise loud instead of returning None into the caller's
+        # non-force_path_c silent `return None` degraded path.
+        raise RuntimeError(
+            f"_cuda_eager.fp8_sparse_mla_apply_cuda_eager: TileLang-CUDA "
+            f"prepared-buffer FP8 Sparse-MLA apply kernel launch/writeback "
+            f"failed ({type(exc).__name__}: {exc})."
+        ) from exc
 
     out = _torch_cuda_to_mlx(
         out_t.reshape(batch, seq_len, heads, d_v), mx_out_dtype
@@ -1304,8 +1342,17 @@ def m2rnn_mapped_packed_post_fwd_cuda_eager(
         )
         kernel(ci_t, W_t, xf_t, h0_t, D_t, proj_t, h_last_t, tanh_t, post_t)
         torch.cuda.synchronize()
-    except Exception:
-        return None
+    except Exception as exc:
+        # RULE #1: a kernel launch/writeback failure is a real bug in the
+        # vendored m2rnn mapped-packed-post fwd kernel, not "feature absent".
+        # Raise loud instead of returning None (the caller would otherwise
+        # re-raise a generic "dispatch failed" that hides the real error, or
+        # fall back to the pure-MLX reference scan).
+        raise RuntimeError(
+            f"_cuda_eager.m2rnn_mapped_packed_post_fwd_cuda_eager: TileLang-CUDA "
+            f"m2rnn mapped-packed-post fwd kernel launch/writeback failed "
+            f"({type(exc).__name__}: {exc})."
+        ) from exc
 
     post = _torch_cuda_to_mlx(post_t, out_dtype)
     h_last = _torch_cuda_to_mlx(h_last_t, out_dtype)

cppmega_mlx/nn/_tilelang/_engine_dispatch.py

@@ -360,16 +360,23 @@ def _engine_with_msl_extraction(
 
     lowering = extract_msl_from_engine_artifact(artifact, target=target)
     if lowering is None:
-        if not _MSL_EXTRACTION_FALLBACK_WARNED:
-            warnings.warn(
-                "cppmega_mlx._tilelang: engine_with_msl_extraction returned "
-                "None (non-metal target or artifact had no kernel_source); "
-                "falling back to MSL shim.",
-                UserWarning,
-                stacklevel=2,
-            )
-            _MSL_EXTRACTION_FALLBACK_WARNED = True
-        return _shim_lower(prim_func, target, pass_configs=pass_configs)
+        # RULE #1: the engine lowering SUCCEEDED (we are past the ImportError
+        # guard) but MSL extraction returned None. That is NOT "engine absent" —
+        # it is extraction silently failing (non-metal target, or a metal
+        # artifact with no kernel_source). Falling back to the legacy MSL-string
+        # shim here is exactly the forbidden "MSL-instead-of-tvm-ffi" silent
+        # fallback: it papers over a real extraction bug with a different
+        # (legacy) lowering. Raise with where+what; if the caller genuinely
+        # wants the shim they must route explicitly via CPPMEGA_MLX_TILELANG_ENGINE=shim.
+        raise RuntimeError(
+            f"_engine_with_msl_extraction: TileLang engine lowering for "
+            f"target={target!r} succeeded but "
+            f"extract_msl_from_engine_artifact returned None (non-metal target, "
+            f"or the engine artifact exposed no kernel_source). Refusing to "
+            f"silently fall back to the legacy MSL shim (RULE #1) — this points "
+            f"at a real MSL-extraction bug for this artifact. Set "
+            f"{_FLAG_ENV}=shim to explicitly use the legacy MSL lowering."
+        )
     return lowering
 
 

cppmega_mlx/nn/_tilelang/m2rnn_path_c.py

@@ -3430,8 +3430,17 @@ def _m2rnn_fwd_path_c_full(
         )
     try:
         kernel, lowering = _fwd_kernel_for(batch, seq, heads, k_dim, v_dim, carrier_dtype)
-    except Exception:
-        return None
+    except Exception as exc:
+        # RULE #1: a compile failure of the selected M2RNN Path C fwd kernel is a
+        # real bug, not a "Path C not applicable" signal (those are the explicit
+        # None returns above for dtype/can_run_metal/h0). Availability is probed
+        # separately by `_kernel_lowering_status`; here, raise with where+what.
+        raise RuntimeError(
+            f"_m2rnn_fwd_path_c_full: M2RNN Path C fwd kernel compile/lowering "
+            f"failed for batch={batch} seq={seq} heads={heads} k_dim={k_dim} "
+            f"v_dim={v_dim} carrier={carrier_dtype} "
+            f"({type(exc).__name__}: {exc})."
+        ) from exc
 
     del lowering
     if out is None:
@@ -3537,8 +3546,15 @@ def _m2rnn_bwd_path_c_kernel(
         )
     try:
         kernel, lowering = _bwd_kernel_for(batch, seq, heads, k_dim, v_dim, carrier_dtype)
-    except Exception:
-        return None
+    except Exception as exc:
+        # RULE #1: a compile failure of the selected M2RNN Path C bwd kernel is a
+        # real bug, not a "Path C not applicable" signal. Raise with where+what.
+        raise RuntimeError(
+            f"_m2rnn_bwd_path_c_kernel: M2RNN Path C bwd kernel compile/lowering "
+            f"failed for batch={batch} seq={seq} heads={heads} k_dim={k_dim} "
+            f"v_dim={v_dim} carrier={carrier_dtype} "
+            f"({type(exc).__name__}: {exc})."
+        ) from exc
 
     del lowering
     dW, dh0, dk, dq, dv, dxf, _scratch = kernel(
@@ -4145,8 +4161,15 @@ def _m2rnn_packed_fwd_path_c_full(
             v_dim,
             carrier_dtype,
         )
-    except Exception:
-        return None
+    except Exception as exc:
+        # RULE #1: compile failure of the selected packed M2RNN Path C fwd
+        # kernel is a real bug, not a "not applicable" signal. Raise where+what.
+        raise RuntimeError(
+            f"_m2rnn_packed_fwd_path_c_full: packed M2RNN Path C fwd kernel "
+            f"compile/lowering failed for batch={batch} seq={seq} heads={heads} "
+            f"k_dim={k_dim} v_dim={v_dim} carrier={carrier_dtype} "
+            f"({type(exc).__name__}: {exc})."
+        ) from exc
 
     del lowering
     if out is None:
@@ -4218,8 +4241,15 @@ def _m2rnn_packed_bwd_path_c_kernel(
             v_dim,
             carrier_dtype,
         )
-    except Exception:
-        return None
+    except Exception as exc:
+        # RULE #1: compile failure of the selected packed M2RNN Path C bwd
+        # kernel is a real bug, not a "not applicable" signal. Raise where+what.
+        raise RuntimeError(
+            f"_m2rnn_packed_bwd_path_c_kernel: packed M2RNN Path C bwd kernel "
+            f"compile/lowering failed for batch={batch} seq={seq} heads={heads} "
+            f"k_dim={k_dim} v_dim={v_dim} carrier={carrier_dtype} "
+            f"({type(exc).__name__}: {exc})."
+        ) from exc
 
     del lowering
     dconv_input, dW, dxf, dh0, _scratch = kernel(
@@ -4332,8 +4362,15 @@ def _m2rnn_mapped_packed_fwd_path_c_full(
             v_dim,
             carrier_dtype,
         )
-    except Exception:
-        return None
+    except Exception as exc:
+        # RULE #1: compile failure of the selected mapped-packed M2RNN Path C
+        # fwd kernel is a real bug, not a "not applicable" signal. Raise.
+        raise RuntimeError(
+            f"mapped-packed M2RNN Path C fwd kernel compile/lowering failed for "
+            f"batch={batch} seq={seq} total_heads={total_heads} k_dim={k_dim} "
+            f"v_dim={v_dim} carrier={carrier_dtype} "
+            f"({type(exc).__name__}: {exc})."
+        ) from exc
 
     del lowering
     if out is None:
@@ -4447,8 +4484,16 @@ def _m2rnn_mapped_packed_bwd_path_c_kernel(
             dy_dtype,
             "float32",
         )
-    except Exception:
-        return None
+    except Exception as exc:
+        # RULE #1: this try wraps shape validation (which raises ValueError on
+        # malformed dy/tanh_cache) and the mapped-packed M2RNN Path C bwd kernel
+        # compile. Both are real bugs — not "not applicable" signals (those are
+        # the explicit `return None` shape checks, which return normally and
+        # never reach this except). Raise with where+what instead of swallowing.
+        raise RuntimeError(
+            f"mapped-packed M2RNN Path C bwd kernel compile/validation failed "
+            f"({type(exc).__name__}: {exc})."
+        ) from exc
 
     del lowering
     outputs: M2RNNPackedBwdOwnerOutputs = (
@@ -4592,8 +4637,15 @@ def _m2rnn_mapped_packed_post_fwd_path_c_full(
             projected_dim,
             carrier_dtype,
         )
-    except Exception:
-        return None
+    except Exception as exc:
+        # RULE #1: compile failure of the selected mapped-packed-post M2RNN
+        # Path C fwd kernel is a real bug, not a "not applicable" signal. Raise.
+        raise RuntimeError(
+            f"mapped-packed-post M2RNN Path C fwd kernel compile/lowering failed "
+            f"for batch={batch} seq={seq} total_heads={total_heads} k_dim={k_dim} "
+            f"v_dim={v_dim} projected_dim={projected_dim} carrier={carrier_dtype} "
+            f"({type(exc).__name__}: {exc})."
+        ) from exc
 
     del lowering
     h_last, tanh_cache, post = kernel(conv_input, W, xf, h0, D, projected)
@@ -4698,8 +4750,14 @@ def _m2rnn_inline_post_bwd_path_c_kernel(
             carrier_dtype,
             "float32",
         )
-    except Exception:
-        return None
+    except Exception as exc:
+        # RULE #1: compile failure of the M2RNN post-residual/gate bwd-from-
+        # recurrence kernel is a real bug, not a "not applicable" signal. Raise.
+        raise RuntimeError(
+            f"M2RNN Path C post-residual/gate bwd-from-recurrence kernel "
+            f"compile/lowering failed for batch={batch} seq={seq} "
+            f"total_heads={total_heads} ({type(exc).__name__}: {exc})."
+        ) from exc
 
     del lowering
     dy_recurrent, dconv_input, dD, dprojected = _m2rnn_post_bwd_owner_outputs(
@@ -4794,8 +4852,14 @@ def _m2rnn_post_residual_gate_fwd_path_c(
             projected_dim,
             carrier_dtype,
         )
-    except Exception:
-        return None
+    except Exception as exc:
+        # RULE #1: compile failure of the M2RNN post-residual/gate fwd kernel is
+        # a real bug, not a "not applicable" signal. Raise with where+what.
+        raise RuntimeError(
+            f"M2RNN Path C post-residual/gate fwd kernel compile/lowering failed "
+            f"for batch={batch} seq={seq} total_heads={total_heads} "
+            f"carrier={carrier_dtype} ({type(exc).__name__}: {exc})."
+        ) from exc
 
     del lowering
     return kernel(y, conv_input, D, projected)
@@ -4874,8 +4938,14 @@ def _m2rnn_post_residual_gate_bwd_path_c_kernel(
             carrier_dtype,
             "float32",
         )
-    except Exception:
-        return None
+    except Exception as exc:
+        # RULE #1: compile failure of the M2RNN post-residual/gate bwd kernel is
+        # a real bug, not a "not applicable" signal. Raise with where+what.
+        raise RuntimeError(
+            f"M2RNN Path C post-residual/gate bwd kernel compile/lowering failed "
+            f"for batch={batch} seq={seq} total_heads={total_heads} "
+            f"carrier={carrier_dtype} ({type(exc).__name__}: {exc})."
+        ) from exc
 
     del lowering
     dy_recurrent, dconv_input, dD, dprojected = _m2rnn_post_bwd_owner_outputs(