[correctness] DCP A2A capture channel adoption: graphs sharing one channel bake host-side staging-slot parity; stale capture stream-key mappings can hand a graph's channel to later captures

[malaiwah]

Found during a performance/correctness review of the v20 r12 release image voipmonitor/vllm:gilded-gnosis-v20-vllmb46c3aa-si35aebc6-fi801d57a-cu132-20260730-r12. Code refs are against the composed r12 tree; all cited sparkinfer files are byte-identical to public commit de04125508c75e748016493f851e6bd4513b515b, so permalinks use that SHA. vLLM caller refs are from the composed r12 vLLM tree (base b46c3aa). Deployment context: GLM-5.2 753B MoE, 4x RTX 6000 Pro (sm120), TP4, DCP2/DCP4, FULL_AND_PIECEWISE CUDA graphs.

This is a proposal-tone hardening report: we could not construct a reachable corruption in the shipped GLM r12 configuration (reasons below), but the safety of the A2A channel under CUDA-graph replay currently rests on two invariants that nothing in the code enforces or asserts. Related memory-side report: #94 (section 1 covers the per-capture-pass channel slabs).

1. All graphs captured in one manager pass share one channel — with host-baked staging-slot parity

Where:

• sparkinfer/comm/pcie/pcie_dcp_a2a.py#L667-L708 (for_stream), #L766-L789 (capture)
• sparkinfer/comm/pcie/pcie_dcp_a2a.cu#L399 and #L460

        if (
            not self.single_channel
            and _is_current_stream_capturing(self.device)
            and self._capture_channel_stack
        ):
            # Independent graph managers may reuse a torch-owned nested stream
            # key. The enclosing capture, not a stale key mapping, determines
            # which IPC channel is safe for this graph.
            channel = self._capture_channel_stack[-1]
            self._channels[key] = channel
            return channel

    // barrier); no host staging memcpys are issued.
    const int slot = slot_++ % 2;

Mechanism: vLLM enters pool.capture() once per cudagraph-manager pass (vllm/v1/worker/gpu_model_runner.py:6857, vllm/v1/worker/gpu/cudagraph_utils.py:464 wrap the entire PIECEWISE+FULL size loop in one graph_capture(...), which enters capture_b12x_dcp_a2a once). So all graphs of a pass (sizes 4..64, FULL and PIECEWISE) deliberately share one channel via the capture-stack adoption above — that is the by-design behavior already quantified in #94.

The subtle part is the staging-slot double buffer. The device barrier itself is replay-safe: self_counter[...] += 1 executes at kernel run time and the flag phase is value % 2 (pcie_dcp_a2a.cu:68-115). But the staging slot is selected on the host (slot_++ % 2) and baked into the captured kernel node's arguments. The protocol is single-barrier by design (.cu:3-8): after the barrier, peers read my staging slot with no tail barrier before my next op's copy phase. Slot alternation is therefore the only thing that prevents op N+1's staging copy on a fast rank from overwriting the slot a slower peer is still reading from op N.

Under graph replay, alternation across the executed op sequence holds only if every graph baked an even number of channel ops. In the shipped GLM config it does, by accident: each MLA layer under DCP decode issues an all_gather_heads + lse_reduce_scatter pair (vllm/model_executor/layers/attention/mla_attention.py:1429, 1545), 2 x 78 = 156 ops per graph, so every graph starts at slot parity 0 and any replay order alternates correctly. Nothing enforces this. A graph that bakes an odd op count (a config where only one of the pair is A2A-eligible over an odd layer count, a partial-layer/breakable graph, a future model) makes two consecutive replays issue back-to-back same-slot ops, re-creating exactly the overwrite window the double buffer exists to close. On-rank replays are stream-serialized, but the hazard window is cross-rank (peer's read phase vs my next copy phase), which stream order does not cover.

2. for_stream during an unwrapped capture can silently adopt another graph's channel

Where: pcie_dcp_a2a.py#L687-L708

        channel = self._channels.get(key)
        if channel is not None:
            return channel
        if _is_current_stream_capturing(self.device):
            ...
            if self._channels:
                channel = next(iter(self._channels.values()))
                self._channels[key] = channel
                return channel

Mechanism: two latent paths, both only reachable when a CUDA graph capture happens without entering pool.capture() (empty _capture_channel_stack):

1. The self._channels.get(key) probe runs before the _is_current_stream_capturing check. Mappings written during earlier captures (lines 685, 696, 700) are never invalidated when the capture context exits, and torch recycles capture-stream handles (the code's own comments at L681-683 and L779-780 acknowledge this). A later unwrapped capture whose stream key collides therefore silently bakes a previous graph's channel.
2. Failing that, next(iter(self._channels.values())) adopts an arbitrary channel. Since _channels is a per-rank local dict keyed by process-local stream pointers, two ranks can in principle adopt different channels for the same logical graph — channels are created collectively in order, so a rank-divergent adoption pairs mismatched IPC slabs, which is a deadlock/corruption, not just a race.

Reachability (honest): in r12 vLLM every A2A-bearing capture is wrapped by graph_capture() -> capture_b12x_dcp_a2a, and the memory-profiling probe uses checkpoint_channels/rollback_channels. We found no unwrapped capture path in the shipped image, so both paths are latent — but they are the API's default behavior for any embedder that captures graphs without the vLLM wrapper.

Suggested fixes

• Derive the staging slot from device-side state instead of the host counter — e.g. from the parity of the existing self_counter (it already increments once per executed op), which makes replay sequences self-consistent regardless of baked arguments and removes the even-op-count invariant entirely. This looks like the smallest complete fix.
• Alternatively: assert at capture exit that the channel's ops-baked count is even (cheap: snapshot slot_ in capture() enter/exit), and document the invariant.
• Invalidate _channels entries created during a capture when the capture context pops, so recycled stream keys cannot resurrect a graph-owned channel (for_stream should reach the "no channel during CUDA graph capture" RuntimeError instead of the stale mapping / arbitrary adoption).

Verification hint

Unit-level, no vLLM needed: build a pool with 2 ranks, capture two graphs each containing an odd number of lse_reduce_scatter ops on the shared capture channel, then replay A, A while delaying one rank's read phase (e.g. stretch the reduction with a spin kernel between barrier and read). The second replay's copy lands in the same slot the delayed peer is still reading; comparing against the eager reference exposes the overwrite. With the device-side slot derivation the same test passes.

[malaiwah]

Independent re-review by a second reviewer (I did not author the original claims), against the same trees — sparkinfer de04125/35aebc6, composed r12 vLLM (base b46c3aa) — with two retractions and one new finding in the lines the report already quotes.

Claim	Verdict	Note
1a. Staging slot is host-side (slot_++ % 2) and baked into captured kernel-node args	HOLDS	Verbatim at .cu:399 and .cu:460; slot_ is used nowhere else, no device-side alternative, no tail barrier in either kernel
1b. The device barrier does not cover the staging slot	HOLDS	start_barrier is a start barrier (pre_sync=true, before the read phase); the write phase of op N+1 is unprotected. The double buffer is genuinely load-bearing — for the intra-layer AG→RS adjacency
1c. AG + RS share one channel → 2 ops/layer	HOLDS	Both call sites resolve to the same pool key; derivation below
1d. Even op count is an accident of GLM having 78 layers	HOLDS-WITH-CORRECTION	Evenness is structural pairing (2×N for any N), not an accident of N. Also true in both shipped configs, not just the gate
1e. Every graph starts at slot parity 0	HOLDS-WITH-CORRECTION	True, but not because captures start from a clean counter — eager warmups inside graph_capture() also advance it
1f. An odd baked count "re-creates exactly the overwrite window the double buffer exists for"	RETRACT as a corruption risk → fragility/hardening only	Same-slot adjacency can only occur at a graph→graph boundary, and every such boundary carries a mandatory TP-group reduction that closes the ≤1-op skew. Not reachable in any TP>1 MLA config, odd count or not
2.1 Stale _channels.get(key) can bake a previous graph's channel	HOLDS (latent, future-facing)	Confirmed unreachable in shipped vLLM; the one unwrapped capture in the tree cannot reach for_stream — named below
2.2 next(iter(...)) can make two ranks adopt different channels → mismatched IPC slabs / deadlock	RETRACT	CPython dict ordering + collective channel creation make the choice rank-identical. The fallback is still wrong, for a different reason
—	NEW	int slot_ overflows to staging_[-1] after 2^31 ops on one channel — reachable in days on the eager channel; details in §5

---

1c — the two ops really do share one channel (worth stating, it's what makes the count 2N)

_get_b12x_dcp_a2a_pool keys on (group, device, total_heads, head_dim, query_head_dim, max_batch_size). From mla_attention.py:1429 and :1545:

• AG: total_heads = local_heads × ws = num_heads × ws, head_dim = output_head_dim = kv_lora_rank (512, since project_before_merge is False whenever dcp_use_a2a), query_head_dim = mqa_q.shape[-1] = 512+64 = 576
• RS: total_heads = cp_attn_out.shape[1] = num_heads × ws, head_dim = 512, query_head_dim = b12x_query_head_dim = kv_lora_rank + qk_rope_head_dim = 576
• both: max_batch_size = min(dcp_max_batch_size, VLLM_DCP_A2A_MAX_TOKENS), identical (dcp_max_batch_size = max_num_batched_tokens, mla_attention.py:832)

Keys identical → one pool → one channel → one slot_. Confirmed.

1d — the invariant is enforced, by the call site, and it is not GLM-layer-count-dependent

Both ops are issued from one place and gated on the same flag:

dcp_use_b12x = self.dcp_b12x and dcp_small_batch          # 1395
...
elif dcp_use_b12x:
    mqa_q = dcp_b12x_all_gather_heads(...)                 # 1429
...
if dcp_use_a2a:
    attn_out = dcp_a2a_lse_reduce(..., use_b12x=dcp_use_b12x, ...)   # 1545/1550

So a graph bakes 2×N ops, even for any N — the "156 because 78 is even" framing understates the guarantee. The residual fragility is narrower than filed and worth stating precisely: the pair can only decouple if the two independent fallback gates disagree, and there is exactly one realistic disagreement — _try_b12x_dcp_lse_reduce bails on cp_attn_lse.dtype != torch.float32 while _try_b12x_dcp_all_gather_heads has no such check. Today B12X_MLA_SPARSE returns fp32 LSE (enforced at b12x_mla_sparse.py:2033), but the NCCL path's own comment (dcp_alltoall.py:967-970, "some MLA backends return it in the activation dtype") shows a bf16 LSE is a live possibility upstream. That, not layer parity, is the one-line change that would produce an unpaired count.

Reachability correction, both shipped configs (the report only claims the gate): the effective launcher default is DCP_A2A_MAX_TOKENS=64 (serve-glm52-v16.sh:21, reached via serve-glm52-v19.sh → v16; the :-16 default is the hybrid launcher only). So every turnkey capture size 4..64 and every gate size ≤6 takes the b12x path and bakes the full 156. PIECEWISE graphs bake zero — attention is in splitting_ops by default (compilation.py:1147), so it runs eagerly on the capture stream during that pass.

1e — the counter is not zero at capture

capture() registers self._channels[key(context.stream)] = channel (pcie_dcp_a2a.py:782) before capture begins, and both runners execute eager warmups inside graph_capture() on that same stream (gpu_model_runner.py:6916-6927; cudagraph_utils.py:477-480). During those warmups the stream is not capturing, so for_stream falls through to _channels.get(key) and returns the graph channel — i.e. warmup ops consume host slot_ at capture time. Parity survives only because warmups also contribute 156 each. Same pairing, not construction.

1f — why I can't build the failing interleaving, and why it isn't buildable

Within a graph the host counter alternates strictly, so same-slot adjacency can only appear at a graph→graph (or replay→replay) boundary. Every such boundary in an MLA+DCP model contains a mandatory cross-rank rendezvous:

• dcp_a2a_lse_reduce returns a head-scattered [B, H/N, D], so recombining the hidden state is impossible without the row-parallel o_proj reduction over the TP group (deepseek_v2.py:1263). This is not incidental — it is arithmetically required by the head scatter.
• vLLM guarantees tp_size % dcp_size == 0 with DCP reusing TP ranks (config/parallel.py:503), so DCP>1 ⇒ tp>1 and the DCP group is always a subset of a TP group.

A one-shot or NCCL all-reduce is a rendezvous: rank A cannot pass it until every peer has arrived, and a peer arrives only after its previous A2A read kernel completed (same-stream program order). So A cannot reach the next A2A op's pre-barrier write phase while any DCP peer is still in the previous op's read phase — exactly the window an odd count would expose. Meanwhile the adjacency the double buffer does protect — AG→RS inside one layer, with no collective between them — is always alternating because the host counter alternates within a graph.

Net: an odd baked count would still be correct in any TP>1 MLA deployment. The suggested unit test (replay A, A with an odd count, stretched read phase) will fail, but only because the synthetic harness has no intervening collective; it demonstrates the raw channel contract, not a reachable serving bug. Recommend adding a second variant with a 2-rank all-reduce between the replays — it should pass, and that pair of tests is the right documentation of the actual invariant.

[INFERENCE] this rests on the TP reduction being a true rendezvous. reduce_results=False / fused-AR variants relocate the collective but cannot remove it, since the head shards must still be summed.

2.1 — reachability, sharpened

for_stream:687 is reachable only when _capture_channel_stack is empty. Every capture site in the r12 vLLM tree that can execute MLA attention runs under graph_capture() → capture_b12x_dcp_a2a: cuda_graph.py:337, breakable_cudagraph.py:182, gpu_ubatch_wrapper.py:287, encoder_cudagraph.py:280, cudagraph_utils.py:515, plus the profiling probe at gpu_model_runner.py:6745 (which additionally uses checkpoint_channels/rollback_channels). The only unwrapped torch.cuda.graph(...) in the tree is v1/spec_decode/gemma4.py:132, which captures masked_emb.get_top_tokens — no DCP collective, and a Gemma4-only drafter absent from GLM-5.2. So the claim's own "latent / API-default" framing is correct and I'd leave it there. Supporting detail: the key is the raw cudaStream_t (pcie_oneshot.py:63-73), so handle recycling is real, not theoretical.

2.2 — retraction

two ranks can in principle adopt different channels for the same logical graph … deadlock/corruption

This does not hold. _channels is a CPython dict, so next(iter(...)) yields the first-inserted entry; channel creation is collective and symmetric (every rank runs the same code in the same order), and capture() overwriting an existing key preserves its position. Every rank therefore selects the same ordinal channel — in practice the eager default-stream channel created by pool.for_stream() at dcp_alltoall.py:134. Divergence would require rank-asymmetric stream usage, which SPMD execution does not produce. The stream keys differ per rank; the order does not, and only order feeds the selection.

The fallback is still wrong, just differently: it can hand a capture-owned channel to a later capture, creating a cross-graph lifetime dependency and shared staging between two independent graphs. That is the defect to fix; the deadlock framing should be dropped.

5 — new: int slot_ overflow indexes staging_[-1]

  int slot_ = 0;                          // :346
  const int slot = slot_++ % 2;           // :399, :460

At INT_MAX the increment is UB; in practice it wraps to INT_MIN. INT_MIN % 2 == 0, but the next op gets -2147483647 % 2 == -1 → staging_[-1], and every other op thereafter. With the member layout (rank_ 0, world_size_ 4, signals_ 8–71, self_signal_ 72–79, staging_[0] 80–143), staging_[-1] is bytes 16–79 — i.e. signals_.signals[1..7] plus self_signal_. The kernel then writes staged attention data straight over peers' barrier-flag arrays: guaranteed hang or silent corruption, not a benign wrap.

Budget: 156 ops per model pass on the eager channel (the graph channel freezes slot_ at capture). 2^31 / 156 ≈ 13.8 M passes — ~1.6 days at 100 eager passes/s, ~8 days at 20/s. Reachable for a long-lived production server that mixes small (≤64-token) non-uniform batches, which dispatch PIECEWISE and therefore run attention — and the A2A — eagerly. Fix is one token: slot_ = (slot_ + 1) & 1, or make it unsigned.

---

Fix-safety notes

F1 — device-side slot derivation from self_counter parity (the report's preferred fix). Sound, and cheaper than it looks.

• Capture-safe: no host state, no allocation, no host sync. Both slot base pointers become kernel params: +64 B of RankStaging, trivially inside the param limit. IPC/VA stability is improved, not risked — both slabs are allocated once at channel construction, so nothing depends on which slot a graph baked.
• Numerics: unchanged. Same bytes, different buffer.
• Cost: one extra load of self_counter[blockIdx.x][0] per block before the staging loop — ≤16 blocks at the default block_limit=16, against ≥200 KiB of staged writes + peer reads per op even at capture size 6. Well under 0.05% of the op's traffic. Two implementation requirements: it must be a non-caching load (ld.relaxed.gpu/volatile, not __ldg), and the value must be broadcast to the whole block, which needs a __syncthreads() before the staging loop — today the first one is the pre_sync inside start_barrier.
• Replay determinism across ranks: parity is per (block, peer). blocks derives from batch/heads/threads/block_limit, identical on every rank for a given op, so block b's counter matches block b's counter on every peer, and writer/reader block pairing is preserved (both stride by gridDim.x * warps_per_block). Counters do diverge between blocks when consecutive ops launch different grids (AG: rows = batch × total_heads; RS: rows = batch × heads_per_rank) — harmless, because the barrier is already block-pairwise (peer_counter[..][blockIdx.x][..]).
• No new invariant: this rides on "every rank executes the identical channel-op sequence with identical grid shapes", which start_barrier's self_counter-driven flag phase already requires for the barrier itself. Strictly not worse than today, and it removes a host-side invariant. Use the pre-increment value for the slot; the flag phase then uses the opposite parity, different buffer, no aliasing.

F2 — assert an even ops-baked count at capture exit. I'd drop this. It needs a new binding (slot_ has no Python-side accessor today — nothing in pcie_dcp_a2a.py reads it), and per §1f it asserts a condition that is sufficient but not necessary: it would hard-fail configurations that are in fact safe, including any channel legitimately carrying an unpaired count (e.g. an RS-only channel, which flash_attn.py:801 / flashinfer.py:1659 would produce if they ever passed use_b12x=True). A comment at .cu:399 plus the pairing note at mla_attention.py:1428-1553 documents it without the failure mode.

F3 — invalidate _channels entries when the capture pops. Lifetime-safe: _new_channel also appends to _all_channels (pcie_dcp_a2a.py:622), so dropping the mapping does not free IPC slabs a replayed graph still points at — worth stating explicitly, because dropping the last reference would unmap peer memory under a live graph. But making for_stream then reach the RuntimeError requires deleting the next(iter(...)) fallback too, which breaks standalone embedders that warm one eager channel and capture without pool.capture() — the documented default path. Safer shape: keep the fallback, tag each channel with the capture that created it, and refuse to adopt a capture-owned channel.

F4 — the obviously-correct alternative (tail barrier, drop the double buffer) would also recover ~half the channel's staging (relevant to #94), but adds a second system-scope barrier per op to a protocol that is one-barrier by design (.cu:3-8). Only worth it if measurement shows the barrier is cheap relative to the op.

What to measure. (a) A2A op latency before/after F1 via the existing SPARKINFER_PCIE_DCP_THREADS / _BLOCK_LIMIT sweep — expect sub-µs, no serving A/B needed. (b) The odd-count unit test in the report plus the with-all-reduce variant; the second one passing is the evidence for the §1f downgrade. (c) Nothing for the overflow fix — it is unobservable until it isn't.

Priority triage

• Trivial + safe, land now: slot_ overflow (& 1); document the AG/RS pairing at mla_attention.py:1428-1553 and the slot invariant at .cu:399/:460; make the next(iter(...)) fallback refuse capture-owned channels.
• Needs measurement: F1 device-side derivation — correctness-neutral hardening that removes the invariant; measure the added load, expect noise.
• Drop: the even-count assert (asserts a non-necessary condition, can reject safe configs); the rank-divergence deadlock claim (retracted, §2.2); the "one accident away from corruption" framing for odd counts (retracted, §1f — the mandatory TP reduction closes the window structurally).

Overall: the two mechanisms are real and correctly located, but the severity is lower than filed — this is hardening plus one genuine (if slow-fuse) overflow bug, not a latent corruption.

[malaiwah]

Implemented the surviving capture-channel ownership fix as PR #103: #103

Scope follows the second-review correction already recorded here: the original odd-operation corruption claim is retracted, and the signed-slot overflow is handled separately by PR #101 / issue #97. PR #103 removes every stream-key alias to a graph-owned channel when pool.capture() exits, restores an eager mapping when capture temporarily replaced it, and retains the graph channel itself in _all_channels for replay lifetime.

Verification:

• CPU-isolated r12 DCP A2A module: 11 passed
• regressions cover target/draft alias retirement, graph lifetime, recycled keys, eager mapping restoration, nested/repeated capture, rollback, and close
• Ruff lint: pass
• independent adversarial review: APPROVE (0.97), no findings
• CodeRabbit: pass

No live-GPU result is claimed. The active production container was not rebuilt or restarted and remains healthy.

[malaiwah]

Completion update: PR #103 now contains the full surviving correction and the reconciled #101 stack: #103

• Capture-context exit restores overwritten eager mappings and removes every alias to the graph-owned channel; _all_channels retains the graph channel for replay lifetime.
• DCP kernels now receive both slabs and derive slot parity from the pre-increment device execution counter, so adjacent same-operation graph replays advance slots by execution rather than capture-time host state.
• Rank-specific AG/RS grid/counter pairing is preserved; shared-memory pointer broadcast occurs before staging.
• CPU ownership suite: 11 passed; bounded-selector regression: 1 passed; DCP/one-shot/two-shot extensions all compiled and loaded; Ruff passes.
• Independent reviews: channel lifetime APPROVE (0.97), device parity APPROVE (0.99).

The original odd-operation corruption claim remains retracted. The retained fixes are capture-channel lifetime and execution-owned replay parity. The real CUDA graph regression is committed but not run because all four GPUs remain occupied by the unchanged production service; no live-GPU result is claimed. Merge #101 before #103.

[malaiwah]

Final implementation update:

• PR #103 head 5c4f8b0 retains graph channels for replay lifetime, restores/removes nested capture stream aliases correctly, and selects DCP slabs from device execution parity.
• The stack includes fix(pcie): prevent two-slot selector overflow #101's zero-extra-launch reusable-collective replay fix and removes the obsolete DCP host selector after migration.

Verification on the current stack: 44 passed, 2 opt-in GPU skips; DCP A2A, one-shot, and two-shot CUDA extensions compiled+loaded together from a fresh cache; Ruff passes. Live multi-GPU execution is assigned to the separate field-test agent.

[malaiwah]

Rental revalidation correction — FAIL pending safe replacement

The earlier PASS at PR #103 head 8fd90a4f687895e705bdf52fcff84ef653a5abf0 proved capture-channel alias restoration and the corrected scratch-reuse numerical oracle, but a later strict synchronization/lifecycle audit found P1 defects outside that oracle. Per-block DCP parity fails a skewed large→small→large grid sequence; DCP teardown is non-retryable/best-effort; partial setup and GC can release live IPC mappings; subgroup ordering is wrong for nonmonotonic groups; and fixed resident-grid barriers are not generically safe.

I corrected the PR validation comment with exact heads, retained positive GPU gates, blockers, and immutable logs: #103 (comment)

A replacement with one operation-wide DCP control node, transactional setup, coordinated retryable teardown, safe GC retention and exact group order is in fresh-cache DCP2/DCP4 testing. This issue remains open pending that result and independent review.

[malaiwah]

Replacement qualification update: the frozen composed successor
31fa6a48116471ce423f0338047453ec1032c202
passes the live DCP4 gates relevant to this issue: canonical pre-preparation
followed by opposite rank-local capture order (1,025 replays), DCP eager/graph
64 with injected retryable teardown, asymmetric/preflight/reduced-SM rejection
before IPC allocation, and the paired 64-node vLLM integration suite. No Xid,
hang, stale handle, capture-time allocation, rank-order mismatch, or ownership
loss was observed.

This does not retroactively pass public #103 head 8fd90a4; it supersedes
the rejected/unqualified repair lineage through 5b4d0d4 and is published as
draft successor #105. Exact commands, results, host/toolchain, harness-mistake
disclosure, and raw-log hashes are in the #103 evidence comment
and the exact-head #105 validation.

I recommend keeping this issue open until #105 is reviewed and merged.

[malaiwah]

Superseding field result: the frozen semantic-channel successor is not yet
deployable.

The focused four-rank eager/capture/torture and atomic vLLM/SparkInfer suites
all passed, but the first full GLM-5.2 startup gate found a contract the focused
tests missed. During determine_available_memory()'s uncaptured warm-up, a
compiled TP all-reduce requested vllm:eager:allreduce on a physical CUDA
stream that was already owned by the active vllm:target:profile semantic
scope. Its descriptor warm-ups run before CUDA itself reports the stream as
capturing, so SparkInfer did not apply its active-scope override, selected the
static eager channel, and correctly refused to alias the two channels:

RuntimeError: CUDA stream key ... is already bound to another logical PCIe oneshot channel

The engine failed before KV allocation/health, so there is no performance or
deployability claim for SparkInfer 31fa6a4 + vLLM be1e289. The SHA-only
appliance image containing them will not be promoted.

The complete call trace makes the repair direction deliberately narrow:

1. while a semantic capture scope is active, route a pre-capture warm-up on
   that scope's owner stream to its top logical channel;
2. do not extend that routing to unrelated side streams, and retain hard
   failures for genuine out-of-scope stream/channel mismatches;
3. apply the rule symmetrically to one-shot and DCP pools; and
4. make failure-path teardown idempotent when Torch has already removed the
   process group from the world-group map.

Re-qualification will include the same full-model boot/traffic/shutdown gate,
not just focused tests.

Complete immutable server log and analysis:

• full rejected log
• campaign section L

[malaiwah]

Resolution candidate is now complete: SparkInfer #105 at bc62980 plus exact field vLLM f99e1e7b8636ca3811ab6d23084ac6da63420dc3 passed the full GLM-5.2 graph/traffic/shutdown gate; clean upstream caller PR local-inference-lab/vllm#216 is patch-identical. The successor preserves canonical collective preparation, semantic owner isolation, opposite rank-local capture order, execution-owned replay state, retryable IPC teardown, and the pre-capture owner-stream rule found by the first failed full boot. Exact evidence: https://github.com/malaiwah/glm52-exl3-vast/blob/5c76a2536e7fc9a5f1cb6bf182531889f5385e65/docs/field-review-results/2026-07-30-vast-46335896/UPSTREAM-REPAIR-CAMPAIGN.md Reproduction: https://github.com/malaiwah/glm52-exl3-vast/blob/5c76a2536e7fc9a5f1cb6bf182531889f5385e65/docs/field-review-results/2026-07-30-vast-46335896/COUNTER-VALIDATION.md. Keeping this issue open until #105 merges.