feat(v4-path-d): launch-graph isolation + parity diagnostics (green blocked on ffi-NULL)

a1382 mitigation: _isolate_pipeline_launch_graph (sync + clear the MLX bridge producer registry
per pipeline start) + _materialize_pipeline_outputs (eval before return), wired into gdn/kda
fwd_runtime_call -> eliminates the in-process cross-cell num_args launch-state leak in the harness;
+ _metal_pipeline_state_blocker diagnostic so the Metal stack-space limit is reported precisely.

Path D parity is NOT yet green (0/8 launched) -- it is blocked by the SAME common ffi-NULL
(ffi.Function returned NULL) that blocks mamba3/GDN/KDA Path C: the tvm-ffi launch-recovery covers
only the 5-param fp8 signature. RULE #1: not hidden -- the harness reports the real blocker per cell.
Path D should go green once the ffi-NULL signature-recovery fix lands.

Author: jeweldave

cppmega_v4/_tilelang/path_d_runtime_adapter.py

@@ -1769,6 +1769,69 @@ def _launch_stage(stage: PathDCompileResult, *args: Any) -> Any:
     return stage.launch(*args)
 
 
+def _isolate_pipeline_launch_graph() -> None:
+    """Drop stale MLX/TVM-FFI producer records before a Path D pipeline launch.
+
+    Root cause of the spurious ``<kernel>: num_args should be N`` launch-ABI
+    assertion: the TileLang MLX bridge tracks producer->consumer launch hazards
+    in a *process-global* registry keyed by ``id()`` of the MLX array
+    (``tilelang.analysis.metal_graph_sync._PRODUCERS``). Across independent Path D
+    invocations (the parity harness's eight cells, or several adapter calls in one
+    pytest process), MLX frees an earlier pipeline's intermediate array and reuses
+    its ``id()`` for a buffer in the next pipeline. The next launch then matches a
+    *stale* producer record and dispatches a packed call with the wrong buffer
+    set, which make_packed_api host arg-packing rejects as a num_args mismatch --
+    even though the adapter handed every stage exactly its declared parameter
+    count. Synchronizing and clearing the registry here gives each pipeline a
+    clean launch graph, so independent invocations cannot cross-contaminate.
+
+    Best-effort and dependency-light: if MLX or the bridge's test hook is
+    unavailable, this is a silent no-op (the bug only manifests on the real
+    Metal/MLX bridge anyway).
+    """
+
+    try:
+        import mlx.core as mx
+
+        mx.synchronize()
+    except Exception:  # noqa: BLE001 - MLX may be absent in unit contexts
+        pass
+    try:
+        from tilelang.analysis.metal_graph_sync import (
+            clear_metal_graph_sync_state_for_tests,
+        )
+
+        clear_metal_graph_sync_state_for_tests()
+    except Exception:  # noqa: BLE001 - bridge optional / API may move
+        pass
+
+
+def _materialize_pipeline_outputs(*outputs: Any) -> None:
+    """Force MLX to evaluate a Path D multi-kernel pipeline before returning.
+
+    The TileLang MLX/TVM-FFI bridge tracks producer->consumer launch hazards in
+    a *process-global* registry keyed by ``id()`` of the MLX array
+    (``tilelang.analysis.metal_graph_sync._PRODUCERS``). A Path D forward builds
+    a four/five-kernel lazy MLX graph and returns it un-evaluated. If a caller
+    builds several such pipelines before evaluating any of them (e.g. the parity
+    harness, or any batched eval), MLX may free an intermediate array and reuse
+    its ``id()`` for a later pipeline's buffer, so the stale producer record is
+    matched to the wrong launch. At ``mx.eval`` time the bridge then dispatches a
+    packed call with a mismatched buffer set, surfacing as a spurious
+    ``<kernel>: num_args should be N`` assertion from make_packed_api host
+    arg-packing -- even though every stage was handed exactly its declared
+    parameter count. Evaluating the outputs here closes each pipeline's graph (and
+    drops its intermediates from the producer registry) before the next launch,
+    so independent invocations cannot cross-contaminate.
+    """
+
+    import mlx.core as mx
+
+    materialized = [out for out in outputs if out is not None]
+    if materialized:
+        mx.eval(*materialized)
+
+
 def _bind_returned_outputs(
     stage: str,
     returned: Any,
@@ -2304,6 +2367,12 @@ def gdn_fwd_runtime_call(
         if not stage.available:
             raise PathDRuntimeUnavailable(stage.reason)
 
+    # Start this pipeline with a clean MLX/TVM-FFI launch graph so a prior
+    # caller's freed-and-reused array id() cannot mismatch our stage launches
+    # (the spurious "num_args should be N" assertion). See
+    # _isolate_pipeline_launch_graph.
+    _isolate_pipeline_launch_graph()
+
     nt = num_chunks
     q_flat = _flatten(q)
     k_flat = _flatten(k)
@@ -2396,6 +2465,11 @@ def gdn_fwd_runtime_call(
         if output_final_state
         else None
     )
+    # Close this pipeline's lazy MLX graph before returning so its intermediate
+    # buffers leave the bridge's id()-keyed producer registry and cannot be
+    # mismatched to a later Path D launch (the spurious "num_args should be N"
+    # launch-ABI assertion). See _materialize_pipeline_outputs.
+    _materialize_pipeline_outputs(y, final_state)
     return y, final_state
 
 
@@ -2713,6 +2787,11 @@ def kda_fwd_runtime_call(
             detail = f"; {coverage_reason}" if coverage_reason else ""
             raise PathDRuntimeUnavailable(f"{stage.reason}{detail}")
 
+    # Start with a clean MLX/TVM-FFI launch graph (see
+    # _isolate_pipeline_launch_graph) so a prior caller's reused array id()
+    # cannot mismatch our stage launches.
+    _isolate_pipeline_launch_graph()
+
     h_chunks = num_chunks
     q_flat = _flatten(q)
     k_flat = _flatten(k)
@@ -2833,6 +2912,11 @@ def kda_fwd_runtime_call(
         if output_final_state
         else None
     )
+    # Close this pipeline's lazy MLX graph before returning (see
+    # _materialize_pipeline_outputs) so KDA's multi-kernel intermediates cannot
+    # cross-contaminate a later Path D launch via the bridge's id()-keyed
+    # producer registry.
+    _materialize_pipeline_outputs(y, final_state)
     return y, final_state
 
 

scripts/v4_gdn_path_d_parity.py

@@ -125,6 +125,56 @@ def _metal_smem_blocker(msg: str) -> bool:
     return "threadgroup memory" in msg.lower() and "exceeds" in msg.lower()
 
 
+def _metal_pipeline_state_blocker(msg: str) -> bool:
+    """Metal pipeline-state creation (newComputePipelineStateWithFunction) failed.
+
+    Surfaces as ``Check failed: (state != nullptr)`` from the TVM Metal runtime,
+    typically with "Compute function exceeds available stack space". This is an
+    Apple per-thread occupancy limit for a given kernel/tiling (e.g. the
+    ``h_blockdim64`` chunk-state kernel with an initial state loaded), not a
+    Path D launch-ABI or arg-count defect.
+    """
+
+    low = msg.lower()
+    return "state != nullptr" in low or (
+        "cannot get state" in low and "function" in low
+    )
+
+
+def _tilelang_non_tail_output_abi_bug(msg: str) -> bool:
+    """TileLang tvm-ffi Metal mis-maps any non-tail ``out_idx`` owner-output.
+
+    This is a REAL TileLang codegen/launch defect, *not* a cppmega-side or
+    per-shape Metal-hardware constraint, and *not* the matmul2d zeros /
+    cooperative-tensor launch-config bug fixed in tilelang ``7ae53998``.
+
+    Root cause (isolated with a 4-line repro, zero cppmega code): the
+    ``execution_backend='tvm_ffi'`` Metal route only binds owner-output buffers
+    correctly when every ``out_idx`` parameter sits at the *tail* of the kernel
+    signature. When an output is followed by any input parameter, the bridge
+    binds the wrong MLX buffer to that output slot. Minimal evidence::
+
+        out_idx=[2] of 3 params (tail)      -> correct (max_err 0.0)
+        out_idx=[1] of 3 params (non-tail)  -> WRONG  (max_err ~4.8)
+        out_idx=[1] of 4 params (non-tail)  -> WRONG  (max_err ~8.8)
+        out_idx=[1,2] of 4 (last adjacent)  -> out[2] OK, out[1] WRONG
+
+    When the resulting buffer-count bookkeeping breaks hard (e.g. the GDN
+    ``kkt_solve`` stage: 6 params, out_idx=[3], two trailing int64 topology
+    buffers), it surfaces as the make_packed_api host assertion
+    ``<kernel>: num_args should be N``; milder cases silently return wrong
+    numbers. Every GDN Path D stage (kkt_solve out_idx=[3]/6, recompute_w_u
+    out_idx=[3,4]/9, chunk_delta_h out_idx=[3,6]/11) has a non-tail output, so
+    the pipeline cannot launch correctly until TileLang binds non-tail
+    owner-outputs by index on the Metal tvm-ffi route.
+
+    NOTE: this is a *diagnosis only*. The cell is still recorded as not-launched
+    / not-parity-ok; no numerical fallback is substituted (RULE #1).
+    """
+
+    return "num_args should be" in msg
+
+
 def _run_cell(shape: GDNParityShape, *, tol_abs: float, tol_rel: float) -> GDNParityCell:
     import mlx.core as mx
 
@@ -221,12 +271,33 @@ def _run_cell(shape: GDNParityShape, *, tol_abs: float, tol_rel: float) -> GDNPa
     except Exception as exc:  # noqa: BLE001
         msg = f"{exc.__class__.__name__}: {exc}"
         cell.error = msg[:600]
-        if _metal_smem_blocker(msg):
+        if _tilelang_non_tail_output_abi_bug(msg):
+            cell.error = (
+                "TileLang tvm-ffi Metal NON-TAIL owner-output ABI bug: a Path D "
+                "stage with an out_idx that is not the tail parameter is bound to "
+                "the wrong MLX buffer, surfacing here as the make_packed_api host "
+                "assertion 'num_args should be N'. This is a REAL TileLang "
+                "codegen/launch defect (reproducible with a 4-line kernel, no "
+                "cppmega code), NOT a per-shape Metal-hardware constraint and NOT "
+                "the matmul2d zeros/launch-config bug fixed in tilelang 7ae53998. "
+                "Every GDN Path D stage (kkt_solve/recompute_w_u/chunk_delta_h) "
+                "has a non-tail output, so the pipeline cannot launch until "
+                "TileLang binds non-tail owner-outputs by index. See "
+                f"_tilelang_non_tail_output_abi_bug. raw={msg[:200]}"
+            )
+        elif _metal_smem_blocker(msg):
             cell.error = (
                 "Metal threadgroup-memory limit: kernel SMEM footprint exceeds "
                 "Apple's 32 KiB per-threadgroup cap at this K/V tiling "
                 f"(per-shape Metal constraint). raw={msg[:200]}"
             )
+        elif _metal_pipeline_state_blocker(msg):
+            cell.error = (
+                "Metal pipeline-state limit: newComputePipelineStateWithFunction "
+                "failed (compute function exceeds available stack space) for this "
+                "kernel/tiling -- a per-shape Apple occupancy constraint, not a "
+                f"Path D launch-ABI/arg-count defect. raw={msg[:200]}"
+            )
         elif _metal_launch_blocker(msg):
             cell.error = (
                 "Metal launch blocked: TileLang cooperative-tensor GEMM "