[Bug] v0.27.0 engine permanently stalls after ~1 min idle on 4-node TP=4 (GB10/sm_121, aarch64): shm_broadcast writer starves, requests never reach scheduler

[joesinvestments]

Bug

On 4-node TP=4 (GB10/sm_121a, aarch64), a v0.27.0 engine serves its first requests normally, then permanently stalls after roughly one minute of engine idleness. Requests after the gap never enter the scheduler (num_requests_running stays 0), the API keeps answering /v1/models and /tokenize, and the head's EngineCore logs repeat:

[shm_broadcast.py:802] No available shared memory broadcast block found in 60 seconds.
This typically happens when some processes are hanging or doing some time-consuming work

every 60s indefinitely. All four rank containers stay alive; all TP workers' last log lines are healthy post-warmup (graph capture complete). Only a full fleet restart recovers. Reproduced 3/3 boots. Keeping the API warm (/v1/models polling every 45s) does NOT prevent it; the stall follows engine-level idleness, not connection idleness.

Sequence, timestamped from one boot: warmup collectives complete across all ranks (so the broadcast queue works) → first 1-2 completions succeed (~1.3s each) → ~1-2 min idle → all subsequent requests hang forever.

Environment

Official vllm/vllm-openai:v0.27.0-aarch64 + two patches required to get this far on sm_121: DeepGEMM repinned to 2fd67329 (#51758) and the masked_mha_available fix (#51920). GLM-5.2 (GlmMoeDsaForCausalLM), TP=4 via --nnodes 4 --node-rank N, mp executor, FLASHINFER_MLA_SPARSE_SM120 + FLASH_ATTN MLA prefill, fp8_ds_mla KV, MTP k=2 spec decode, dual-rail RoCEv2, VLLM_USE_BREAKABLE_CUDAGRAPH=0 set. torch 2.13.0+cu130 bundled NCCL.

Related: #51758 validated v0.27.0 on 2x Spark (TP=2 and PP=2) with a 400-request soak and zero hangs — consistent with our result that the first requests DO work; the failure needs the multi-node idle gap, and possibly the 4-node width. This looks adjacent to the pre-0.27 TP-hang family that release fixed for 2-node.

Full per-rank logs, exact launch argv, and timings available; the cluster reproduces this in ~20 minutes per attempt and I can run instrumented builds or targeted env matrices on request.

[ThinkCode]

We have not reproduced this ourselves yet — the below is from reading shm_broadcast.py at v0.27.0, so treat it as a hypothesis rather than a finding. We do run the same shape (4× DGX Spark GB10 / sm_121a, aarch64, TP=4 via --nnodes 4, mp executor) and can run instrumented builds if that's useful.

The obvious explanation is ruled out by the code

The natural first guess is a lost wakeup: readers busy-spin, then after busy_loop_s drop into SpinCondition.wait() → poller.poll() on a SUB socket with CONFLATE=1, while the writer notifies over PUB with SNDHWM=1. PUB/SUB drops silently, so a notification published during the busy→idle transition could be missed.

That can't produce a permanent stall, because the poll is always bounded:

def timeout_ms(self) -> int:
    wait_ms = SHM_READER_RECHECK_INTERVAL_MS          # 5000
    if self.warning_wait_time_ms is not None:
        wait_ms = min(wait_ms, self.warning_wait_time_ms)
    if self.timeout is None:
        return wait_ms                                 # <-- never None

With SHM_READER_RECHECK_INTERVAL_MS = 5000, a reader waiting indefinitely still wakes at least every 5 s and re-runs check() via the continue in acquire_read. A dropped ZMQ ping would self-heal within 5 s.

So both sides are awake and polling on schedule, and they disagree about the buffer state. That makes this a state-divergence deadlock, not a missed-notification bug — worth saying explicitly, because PUB/SUB is where most people (including me) will look first.

Where the two sides can disagree

Writer, acquire_write:

written_flag = metadata_buffer[0]
return not (written_flag and read_count != self.buffer.n_reader)

blocks while the block at its current_idx is written and read_count != n_reader.

Reader, acquire_read:

read_flag = metadata_buffer[self.local_reader_rank + 1]
written_flag = metadata_buffer[0]
return not (not written_flag or read_flag)

proceeds only when the block at its current_idx is written and its own flag is clear.

A permanent standoff needs one of:

1. current_idx divergence — writer waiting on block N, reader waiting on block M≠N. Each sees a legitimately un-actionable block and neither advances.
2. n_reader over-counting — writer waiting for a read_count that can never be reached because some counted reader never writes a flag into this segment. On 1-GPU-per-node × 4 nodes, it'd be worth confirming what n_reader is set to on the head versus how many readers actually touch that local segment; if remote ranks are counted but served by a different transport, read_count can never reach n_reader. That would also fit 4-node failing while the 2-node soak in [Bug]: upgrade vllm from 0.26.0 to 0.27.0 run deepseek v4 flash error #51758 stayed clean.

The idle correlation fits either: the busy→idle transition is the only place the two sides' timing assumptions change.

One free thing that narrows it immediately

LONG_WAIT_TIME_LOG_MSG is logged from two places — line 682 in acquire_write (writer) and line 803 in acquire_read (reader). Knowing which one your head is emitting splits the hypothesis space in half before anyone instruments anything. If it's 682, the writer is waiting on readers; if 803, a reader is waiting on the writer; if both, you have the standoff directly.

What we'd dump to confirm

At stall, per rank: current_idx, the full metadata_buffer bytes for that index (written flag + per-reader flags), local_reader_rank, and buffer.n_reader vs the count of readers actually attached to that segment. If the two current_idx values differ, it's (1); if they match and read_count is stuck below n_reader with no reader claiming that slot, it's (2).

Happy to build an instrumented image and run it on our 4-node fleet — reproduction here would take a couple of hours to stand up, and we can do a targeted env matrix (VLLM_RINGBUFFER_WARNING_INTERVAL, busy-loop duration) alongside it if that helps.

[SyaOtiLan]

This may be closely related to #51593 and the final two fixes in #51538 (47f6574 and db39e67).

Both cases involve the SM120 sparse-MLA backend with MTP (next_n > 1) and a permanent worker stall after the active batch drains, followed by the same shm_broadcast timeout. In #51593, padded MTP slots produced negative context lengths, causing the sparse top-k kernel to deadlock.

Could you A/B test:

1. the same configuration with MTP/speculative decoding disabled; and
2. the current configuration with those two commits from [Bugfix] Make DSV4 sparse MLA work end-to-end for plain decode, MTP, and DSpark #51538 applied?

If either removes the idle-transition stall, that would distinguish the existing MTP drain/kernel deadlock from a separate TP=4 transport or executor issue.

[joesinvestments]

Both replies are extremely useful. Answering the free diagnostic first, since I had the data sitting in my logs, then the A/B plan and some detail that I think favors the MTP hypothesis.

The diagnostic: it is the READER

@ThinkCode's split is decisive and costs nothing, so here it is. The warning our head emits is:

(EngineCore pid=280) INFO 08-12 03:34:43 [shm_broadcast.py:802]
No available shared memory broadcast block found in 60 seconds.

repeating every 60 s indefinitely. In the image (path is vllm/distributed/device_communicators/shm_broadcast.py, same file, same line numbering as quoted above):

649:    def acquire_write(self, timeout: float | None = None):
682:                            LONG_WAIT_TIME_LOG_MSG, VLLM_RINGBUFFER_WARNING_INTERVAL
762:    def acquire_read(
803:                            LONG_WAIT_TIME_LOG_MSG, VLLM_RINGBUFFER_WARNING_INTERVAL

Our line 802 is the logger.info( call whose message argument is on 803, i.e. the acquire_read site (def at 762), not the writer's at 682. Only the reader warning appears; there is no 682 in any rank's logs.

And the process emitting it is the head's EngineCore, not a worker. So the head is blocked reading, waiting on a block that no writer is producing.

That points away from the two-sided standoff in hypothesis (1)/(2) and toward "the thing that should be writing has stopped", which is what a worker deadlocked inside a kernel would look like from here. It also fits @SyaOtiLan's pointer at #51593 (MTP hang after the batch drains) better than a transport bug: the queue demonstrably works, because warmup collectives complete on all four ranks and the first requests return in ~1.3 s. It stops working only after the batch drains.

Supporting detail, all four rank containers stay alive at the stall, and every worker's last log line is a clean post-warmup entry (graph capture finished, JIT monitor activated). No worker error, no exit, no NCCL warning: they simply go quiet.

Config, for the MTP correlation

GLM-5.2 (GlmMoeDsaForCausalLM, QuantTrio Int4-Int8Mix, num_nextn_predict_layers=1), TP=4 across 4 nodes (--nnodes 4 --node-rank N, mp executor), FLASHINFER_MLA_SPARSE_SM120 decode + FLASH_ATTN MLA prefill, --kv-cache-dtype fp8_ds_mla, --max-model-len 200000, --max-num-seqs 6, --max-num-batched-tokens 8192, MTP num_speculative_tokens=2, quantization: compressed-tensors, draft_sample_method: probabilistic, VLLM_USE_BREAKABLE_CUDAGRAPH=0, DeepGEMM repinned 2fd67329, plus the #51920 workaround. Dual-rail RoCEv2, torch 2.13.0+cu130.

So: MTP on, sparse MLA on SM120, stall immediately after drain. That is three of the ingredients in #51593.

One data point from the other direction

On my previous stack (vLLM ab666069, same model, same MTP-on-sparse-MLA shape, same 4 nodes) I see what may be a milder form of the same thing: the first request after a period of idleness takes 30-35 s, then subsequent requests are 2-3 s. It recovers rather than hanging. Same drain-then-idle trigger, same speculative path, older tree. I have not instrumented it and I am not claiming it is the same bug, but if #51593's mechanism has a "recovers" and a "does not recover" variant depending on tree, that older-tree behavior may be a cheaper reproduction target than a full 4-node 0.27 stand-up.

A/B plan

I will run both of @SyaOtiLan's tests on this fleet and report either way:

1. identical config with speculative decoding disabled entirely;
2. current config plus 47f6574 and db39e67 from [Bugfix] Make DSV4 sparse MLA work end-to-end for plain decode, MTP, and DSpark #51538.

Timing note: this is a production cluster serving a live agent workload and it has already spent most of tonight down for this investigation, so both boots happen tomorrow rather than right now. If #51593's fix lands before I get to (2), I will test the merged version instead.

@ThinkCode also runs 4-node GB10 and offered instrumented builds; between us there are now two clusters of the exact shape that fails, which should make confirmation quick. If either of you wants a specific dump at the stall (per-rank current_idx, metadata_buffer for that index, local_reader_rank, n_reader versus attached readers), name the patch and I will run it with the rest.

[joesinvestments]

Both A/B tests run on the 4-node GB10 fleet. Your hypothesis is confirmed, and #51538's Python half alone is sufficient for this failure mode.

Method

Same cluster, same weights, same launcher, one variable per boot. Test procedure per config: serve a turn, go idle 150 s (well past the ~60 s at which the stall was previously observed), serve another turn, repeat 4 cycles; any timeout or any post-idle turn over 10 s is a failure. Predicate validated against a known-good server first so a pass could not be vacuous. Raw JSONL below.

Results

config	idle survival	post-idle latency
v0.27.0 + MTP k=2 (unpatched)	jams permanently, 3/3 boots	never returns
v0.27.0, speculative decoding disabled	4/4 cycles OK	0.3 s
v0.27.0 + 47f6574 + MTP k=2	4/4 cycles OK	0.3 s

{"label":"v027-nospec","cycle":1,"verdict":"OK","pre_idle_s":0.3,"after_idle_s":0.3,"idle_s":150}
{"label":"v027-nospec","cycle":2,"verdict":"OK","pre_idle_s":0.3,"after_idle_s":0.3,"idle_s":150}
{"label":"v027-nospec","cycle":3,"verdict":"OK","pre_idle_s":0.3,"after_idle_s":0.3,"idle_s":150}
{"label":"v027-nospec","cycle":4,"verdict":"OK","pre_idle_s":0.3,"after_idle_s":0.3,"idle_s":150}
{"label":"v027-51538-mtp","cycle":1,"verdict":"OK","pre_idle_s":1.6,"after_idle_s":0.3,"idle_s":150}
{"label":"v027-51538-mtp","cycle":2,"verdict":"OK","pre_idle_s":0.4,"after_idle_s":0.3,"idle_s":150}
{"label":"v027-51538-mtp","cycle":3,"verdict":"OK","pre_idle_s":0.3,"after_idle_s":0.3,"idle_s":150}
{"label":"v027-51538-mtp","cycle":4,"verdict":"OK","pre_idle_s":0.3,"after_idle_s":0.3,"idle_s":150}

So disabling MTP removes the stall, which points at the speculative path exactly as you suspected, and applying #51538 removes it while keeping MTP enabled. On this hardware the fix resolves #51921.

One practical finding about testing #51538 from a released image

Of the two commits you named, only 47f6574 can be applied to vllm/vllm-openai:v0.27.0-aarch64. It patches vllm/v1/attention/backends/mla/indexer.py, which ships in the wheel. db39e67 patches csrc/libtorch_stable/cooperative_topk.cuh and persistent_topk.cuh, which are CUDA sources compiled into the extension and absent from the installed tree, so it cannot be applied without building vLLM from source.

That means the result above was obtained with the Python half only, and it was sufficient to eliminate the idle stall in 4/4 cycles. If db39e67 fixes a distinct symptom (a wrong result rather than a hang, say), that would not be visible in this test. Worth knowing that anyone validating this PR from the official image is testing half of it by default.

Stack detail, for the record

GLM-5.2 (GlmMoeDsaForCausalLM, QuantTrio Int4-Int8Mix), TP=4 across 4 nodes, mp executor, FLASHINFER_MLA_SPARSE_SM120 decode + FLASH_ATTN MLA prefill, fp8_ds_mla, ctx 200000, max_num_seqs 6, MTP num_speculative_tokens=2 with quantization: compressed-tensors and probabilistic sampling. Two other patches were required to reach a serving state on sm_121 at all: DeepGEMM repinned to 2fd67329 (per #51758) and the masked_mha_available workaround from #51920.

Happy to re-run any of this, extend the idle duration, or test a source build with both commits if that would help close #51538.

[joesinvestments]

Correction to my previous comment: my pass result was premature. Do not treat #51538 as confirmed on this hardware.

After posting, I migrated this config to production and immediately ran a heavier battery against it: a 12-way concurrent storm and a 30K-token cold prefill, rather than the light single turns my idle test used. The engine failed within the hour, and my external monitor recorded WEDGED: idle, generation probe did not return in 90s on the drain that followed.

So the honest statement of the result is narrower than what I posted:

• v0.27.0 + 47f6574 + MTP survived 4/4 idle cycles of 150 s each with light single-turn traffic.
• The same build failed after concurrent and deep-prefill load, on the drain that followed.

My test was too weak to support the conclusion I drew from it. The distinguishing variable appears to be what precedes the drain: light single-stream traffic drains cleanly, while a concurrent storm or a long prefill does not. That is a more specific reproduction recipe than I had before, and I would rather hand you that than a false pass.

Two things this does not change: disabling speculative decoding still removed the stall in 4/4 cycles, so the speculative path remains implicated; and db39e67, the CUDA half of #51538, still could not be applied to the released image, so everything above tests the Python half alone. It is entirely possible the kernel half is what closes this, which would also explain why the Python half alone survives light traffic and not heavy.

Next from me, and I will report either way: re-run with load-then-drain as the explicit trigger (storm, then idle, then probe), several cycles, on both the patched and unpatched builds, so the comparison is like-for-like at the traffic shape that actually breaks it. If a source build with both commits would be more useful to you than more black-box testing, say so and I will build it.

Apologies for the noise on a thread you are using to make decisions.

[SyaOtiLan]

@joesinvestments Thanks for correcting the result — the load-to-drain trigger is more useful than the initial light-idle pass.

This still implicates the speculative path, but it does not yet validate #51538: 47f6574 clamps the known negative lengths in the Python indexer paths, while db39e67 hardens the CUDA top-k kernels against any out-of-range per-row length.

A source build with both commits, tested with the same concurrent/30K-prefill → drain → probe sequence, would be the most decisive next A/B. If possible, running the same heavy sequence once with speculative decoding disabled would also provide a like-for-like control.

If the full pair passes while Python-only fails, that would show the kernel guard is load-bearing on this configuration. If it still fails, we likely have another MTP drain path to instrument.

No apology needed — the correction and stronger reproducer are valuable.

[joesinvestments]

Source build done, both asks run, plus three follow-ups. #51538 does not resolve this on our hardware, and the failure survives disabling speculative decoding entirely — which argues against the MTP-drain mechanism we were all working from.

Build provenance, since the whole point was testing the CUDA half

v0.27.0 (4bdc8a788d) built from source with 47f6574 and db39e67 applied and committed, plus #51395. Verified against pre-registered checks on the resulting image, not on the build log: the clamp expression and its comment present in the installed indexer.py, params.lengths[row] > 0 ? params.lengths[row] : 0 present in cooperative_topk.cuh, compiled extension present, version 0.27.x, SM120 declaration touched. 7/7.

Two practical notes for anyone else validating #51538 this way:

• The PR does not apply to the v0.27.0 tag. It targets newer main and conflicts in flashinfer_cutlass_moe.py, gpu/model_runner.py and others.
• Its commits are not reachable from refs/pull/51538/head (rebased), so git cherry-pick 47f6574 fails with bad revision. Applying the two .patch files directly is what works.

Results

Trigger, identical for every row: C=6 storm of 1.2K-token prompts, then a 27K-token cold prefill, then drain, then probe. Verdict rule: a timeout is only a signal; the harness then asks the engine to serve a small fresh request twice, and only a failure to serve counts as wedged. Boots are readiness-gated so cudagraph warmup cannot be scored as a hang.

configuration	verdict
prebuilt v0.27.0 + 47f6574 only, MTP k=2	wedged, cycle 1
source build, 47f6574 + db39e67, MTP k=2	wedged, cycle 1
source build, same image, speculative decoding removed	wedged, cycle 1
source build, max-num-batched-tokens 8192 to 2048	wedged, cycle 1
source build, gpu-memory-utilization 0.91 to 0.85	wedged, cycle 1

All five die in the storm phase, before the drain. Signature is unchanged: /v1/models and /metrics keep answering 200 while completions stop; head EngineCore logs No available shared memory broadcast block found in 60 seconds from acquire_read every 60 s indefinitely.

For contrast, on the same four nodes, the same weights and the same harness, our pre-0.27 stack (ab666069 lineage) survives a heavier storm (C=12) repeatedly.

What I think this means, stated as narrowly as the data allows

Disabling speculative decoding does not prevent the failure, so on this configuration the deadlock does not require the MTP path. That is a genuine surprise given #51593's mechanism, and it means our #51921 is probably not the same bug as #51593 even though the symptoms overlapped. 47f6574 and db39e67 may still be correct and necessary for the MTP case; they are simply not sufficient for whatever kills this configuration under concurrency.

Two things I have ruled out with single-variable cells, in case they were on your list: activation buffer size and memory headroom.

Correction I owe you

Earlier in this thread I attributed our stalls to the speculative path and reported a light-traffic pass as a validation of #51538. Both were premature. The light-traffic pass was retracted the same day; this comment retires the MTP attribution as well. Our public write-up has been corrected to match.

What I can run next

The cluster is free for this and reproduces in about 25 minutes per configuration. I can run an instrumented build, a specific env matrix, VLLM_USE_V2_MODEL_RUNNER=0, a DCP or backend variation, or dump per-rank current_idx / metadata_buffer / n_reader state at the stall if you tell me where to instrument. Full per-rank logs, the exact launch argv, the harness and the raw JSONL are available on request or in the write-up.

[joesinvestments]

Follow-up since the source-build result above. Two things closed, one new bug found and set aside, and the first live capture of the actual stuck state.

Memory pressure is closed as a necessary condition, with a positive control, not just an absence of evidence. Every wedge up to this point, cudagraph or otherwise, coincided with a fresh driver-level NV_ERR_NO_MEMORY burst, so it was a reasonable thing to chase. It took real digging to close cleanly: --gpu-memory-utilization alone barely moved available memory, because --kv-cache-memory-bytes, once set explicitly, does not respect gmu at all per vLLM's own boot log. The dominant fixed allocation is KV cache sized against --max-model-len, and it has a hard floor (10.19 GiB at max-model-len 200000 on our config). Dropping max-model-len to 64000 (still well above our 20K-token deep-prefill test) got the fleet to 4-8 GiB genuinely free, verified via /proc/meminfo, not inferred. Reproduced the wedge anyway, at production max-num-batched-tokens 8192, NONE cudagraph mode.

All four cudagraph modes wedge, not just eager. FULL, PIECEWISE, NONE, and eager all hit the identical storm-phase-timeout signature.

Both available backend alternatives on this hardware are closed, checked against source before spending a boot cycle, then confirmed live. For device capability 12, vLLM's own platform code lists exactly two MLA candidates: TRITON_MLA and FLASHINFER_MLA_SPARSE_SM120. Forcing TRITON_MLA gets a clean, fast rejection: ['kv_cache_dtype not supported', 'sparse not supported'], and 'sparse not supported' isn't contingent on the other reason, so this model's DSA sparse indexer can't run on it at all. MoE backend is less locked down (Triton WNA16 has no capability gate), but it fails a different way: weight loading completes, then assert weight_quant.strategy == "group" fails constructing the MTP drafter's layer specifically. Marlin is the only MoE backend this checkpoint boots under.

A second, different, reproducible bug turned up and got closed as a separate thing, not this one. At the same healthy-headroom config but with max-num-batched-tokens reduced to 2048 (a setting introduced for our own memory diagnostic, not a production value), the engine doesn't hang, it crashes:

File "vllm/v1/core/sched/scheduler.py", line 1761, in update_from_output
    req_index = model_runner_output.req_id_to_index[req_id]
KeyError: 'chatcmpl-...'

Root cause against installed source: sparse_mla_attention.py's forward_mha falls through to the attention base class's unimplemented default whenever prefill_max_seq_len <= topk_tokens, which chunked prefill under max-num-batched-tokens=2048 can satisfy on every chunk of a long prompt. All four ranks hit the identical exception simultaneously; the scheduler KeyError is a downstream artifact of workers dying before producing output. Re-ran at production's max-num-batched-tokens 8192, same everything else: the crash doesn't happen, the original silent hang does instead. Filed separately as #52105 since it's a real bug but not this one.

Live stacks from the actual stuck state, SYS_PTRACE added to the container, py-spy/cuda-gdb installed, same healthy-headroom production-batched-tokens config, sampled twice several minutes apart with zero movement on any rank:

• EngineCore (rank 0's scheduler): blocked in acquire_read on the shm_broadcast queue, waiting for a worker response that never arrives, exactly matching the recurring log line, now confirmed at the stack level.
• Worker rank 0: stuck inside moe_wna16_marlin_gemm.
• Worker rank 1: stuck inside a different compiled kernel entirely.
• Worker ranks 2 and 3: stuck identically inside FlashInfer's sparse-MLA Triton index-conversion kernel (triton_convert_req_index_to_global_index).

No NCCL frame visible in any Python stack. The four ranks are stuck in different kernels, not the same one, which reads as prior divergence rather than a single shared kernel hang: whichever rank finishes first would block on a collective buried inside a later kernel launch that a stalled sibling never reaches. cuda-gdb attach against the live worker did not complete within several minutes and was killed rather than left running; that the debugger itself couldn't interrupt the CUDA context is its own small data point toward a genuine device-side stall.

Full writeup with exact numbers and the raw RAS/memory snapshots: https://github.com/joesinvestments/GLM-5.2-QuantTrio-TP4-DCP2-4x-DGX-Spark/blob/main/v027/MEMORY-AND-KERNEL-FINDINGS.md

Open question we don't have an answer to yet: what would actually distinguish "prior rank divergence that later deadlocks on a collective" from "each rank independently wedged in device state for unrelated reasons that happen to look similar." Native/CUDA-level frames underneath the four kernels would settle it; the aarch64 unwind path in py-spy's own --native mode failed here (UNW_EBADREG), and cuda-gdb couldn't attach live. Open to suggestions on getting device-side visibility on this platform if anyone's done it on GB10.

[SyaOtiLan]

@joesinvestments Thanks for the thorough follow-up. This rules out memory pressure and CUDA Graph mode as necessary conditions, and #52105 appears to be a separate crash.

One important distinction: forcing TRITON_MLA does not test the sparse Triton fallback from #49026. TRITON_MLA is a dense backend and is expected to reject sparse attention. #49026 instead provides an sm12x-compatible Triton implementation behind FLASHMLA_SPARSE, which is not present in the tested v0.27 tree.

The different Python stack locations also do not identify the original device kernel; they remain compatible with the FlashInfer sparse-MLA device-queue livelock captured in #49026.

@marksunner, since this closely overlaps your #49026 investigation, does this capture look consistent with the same device-queue livelock?

@pavanimajety, as the MLA backend code owner, would porting the sparse Triton fallback from #49026 be the preferred vLLM-side direction here?

I do not have access to an SM121 cluster to validate a current-main port, so I’ll leave the implementation and device-level confirmation to those with the required environment. A future port of #49026 could be validated with the reporter’s existing C=6 workload.

[joesinvestments]

Correction taken, and it's the right one.

TRITON_MLA was not a test of #49026's fix. It's a plain dense backend, unrelated to #49026, and it was simply the only other candidate vLLM's source lists for capability 12 in our current, unpatched tree. #49026's actual proposal, a capability-12-compatible FLASHMLA_SPARSE, isn't present in that tree at all, its own capability gate excludes us on this version, which is a fact I'd confirmed earlier and should have been explicit about instead of letting "closed the attention-backend lever" read as broader than it is. What's actually closed: no attention backend alternative exists in our current tree. What's not closed: whether a capability-12-compatible sparse Triton path, if it existed here, would change the outcome. Different claims, and I conflated them.

Same correction applies to the stack read. Four ranks in four different Python-level call sites is real data, but it doesn't reach device state, so it's consistent with genuine prior divergence and equally consistent with the FlashInfer device-queue livelock in #49026. I stated my own reading with more confidence than the evidence actually supports. Should have said "consistent with" both, not favored one.

If a current-main port of #49026 happens, we're glad to be the validation environment. Same C=6 storm plus 20-27K cold-prefill trigger that's produced every result in this thread, same hardware, and now with a dual RAS/memory snapshot on any wedge so a live capture (or a survive) comes with real data attached either way.

[marksunner]

@SyaOtiLan - direct answer to your question first: yes, this capture is consistent with the #49026 device-queue livelock on every axis the capture can show. Cross-rank stacks in four different kernels is the livelock's expected presentation, not evidence against it - under a jammed device queue, each host thread freezes at whatever cuLaunchKernel it issues next, so the kernels named in py-spy are the jam's victims, not its cause. We learned that the expensive way: we spent two rounds indicting marlin because our hosts were piled in moe_wna16_marlin_gemm (rank 0's exact op here) before cuda-gdb on a live wedge named the real spinner - a flashinfer sparse_mla_sm120-family kernel in an mbarrier phase TRYWAIT loop. The rest matches too: no NCCL frames anywhere (the wedged step's collective is never launched), the head starving in acquire_read (workers behind a jammed queue never respond), spec-decode independence, and the attach-resistant wedge state - we hit one of those as well and logged it as its own small datum. The honest limit: consistency is not confirmation. Python stacks structurally can't name the guilty device kernel, and the device-level receipt on your tree doesn't exist yet. Our validation is one config, one fleet (GB10, our 0.23.1-era tree); plausibly family-wide across sm12x is as far as I'll claim.

@joesinvestments - the falsification table, the "timeout is only a signal" rule, and both corrections in your last comment are exactly how an investigation like this stays trustworthy, and "consistent with both" is the right resting point for what your capture shows. Here's how to split it on your own cluster, cheaply - with one trap flag on the instrumentation you've planned: a plain RAS status query will read healthy during this wedge. In every one of our eight captures, RAS reported comm RUNNING OK, 4/4 ranks while the fleet was fully wedged, so a "dual RAS/memory snapshot" as phrased will come back clean and settle nothing. The discriminator is the per-rank collective op counts, dumped twice, ≥90 s apart (the RAS endpoint listens on 127.0.0.1:28028 inside the rank processes; NCCL ≥2.24):

• counts frozen at pre-request values on all ranks, spread 0–1 ops, no in-flight collective, zero movement between samples → kernel-upstream livelock (the [Bugfix][Platform] sm12x: add FLASHMLA_SPARSE (Triton sparse-MLA) to the MLA-sparse backend candidate list — flashinfer sparse_mla_sm120 livelocks on GB10 #49026 class); the wedged step's collective was never launched. That was the signature of all 8 of our wedges. (Don't read a 1-op spread as divergence - it's a victim-timing artifact of when each rank hit the jam.)
• counts advanced into the request era with a rank parked mid-collective → genuine prior divergence deadlocking on a collective - a different disease, and the attention route would be the wrong lever.

Pair it with nvidia-smi on all four nodes: our livelock read ~96 % GPU util at ~18 W with 0 % memory utilisation, zero movement across samples — a spin loop reports busy while doing no work. Those two captures on your next wedge turn "consistent with both" into a verdict, on your tree, for free. Full receipts: https://github.com/marksunner/glm52-dgx-spark-deadlock-evidence and the results comment on RFC #48720.

On the validation offer - genuinely appreciated, and I'd rather answer it straight than leave it hanging: we're not in a position to put our cluster or engineering time into this at the moment unfortunately, or to commit to bringing the branch up to current main right now. What I can do is share everything we've already banked, and all of it is public. The #49026 reference branch carries the working stack: the Triton sparse-MLA kernels are @jasl's - taken verbatim from the deepseek_v4 path of the jasl/vllm fork, and that credit should travel with any derived port - and the branch history shows exactly how they were wired into the backend selection, the model-side support, and the indexer/MTP paths against a July main. Fair warning for anyone who picks it up: it's a port, not a rebase. The kernels carry over near-verbatim in self-contained new files, but flashmla_sparse.py's build() has been rewritten upstream since our base and the model-side support overlaps upstream's SP-gather changes, so the wiring needs re-derivation against current code, not conflict resolution. One caveat on the route itself: our correctness validation was operational - byte-identical greedy smoke across every boot, coherent long completions across 560+ max-context sessions, throughput at or above baseline - not a logit-level audit; a port to main should get real correctness CI before anyone trusts it past "doesn't wedge." Whether a capability-12 Triton sparse path is a direction vLLM wants is the MLA backend owners' call, not mine.

Cheapest next step regardless: instrument the next wedge and read the RAS op-count delta. Frozen counts means it's this disease on your tree and the reference branch shows a route that survived it; advanced-with-one-parked means the attention route was never the lever, and your snapshot will have pointed at the right one instead of a mispatch. Either way, everything we have on the mechanism is published.

[joesinvestments]

@marksunner thank you, that is exactly the kind of answer that moves an investigation. The discriminator is now wired in, and I want to be precise about what has and has not happened since.

Instrumentation, live on the fleet:

• The recovery path now runs a capture step before any teardown: ncclras -f json -vv (built from src/ras/client.cc at the bundled 2.29.x, run on each host since the ranks are --network host) sampled twice 95 s apart on all four ranks, and the script diffs the per-rank collective_counts per communicator between the two samples, alongside nvidia-smi util / memory-util / power / SM clock, container state, vllm:num_requests_running, and the last 200 log lines per rank, all written to a directory the recovery cannot touch. Verified both branches on the live fleet tonight: idle, moved=0 on all 16 rank x comm pairs (which is why the running-request count is recorded next to it, an idle fleet is frozen by definition); one request in flight, moved=8 per node with a 0-1 op spread across ranks, exactly the victim-timing artifact you described. Verdict rule is yours verbatim: frozen with requests in flight plus high util at low power and 0 % memory util = kernel-upstream livelock; advanced with one rank parked = genuine divergence. Your warning about a plain RAS status query reading healthy is in the script header; the summary view indeed says RUNNING OK while the JSON counters are what carry the signal.
• Separately (for RoCEv2 receiver deadlocks after successful RDMA_WRITE completions land on 'unused' request slots (net_ib/p2p.cc resiliency path) NVIDIA/nccl#2334, the RoCE receiver desync, a different failure) the ranks now log NCCL_DEBUG=INFO, SUBSYS=INIT,NET to a persistent host-mounted file, so a full per-rank log survives the restart.

What has changed on the attention path since my last comment: production moved off FLASHMLA_SPARSE (the FlashInfer sparse_mla_sm120 kernel your cuda-gdb named) to B12X_MLA_SPARSE, the CuTe-DSL sparse-MLA + sparse-indexer path for GB10 that 0xdfi's O14 work is built on. Nothing FlashInfer sits in the sparse attention path any more. That is the same class of move as the #49026 route (a non-FlashInfer sparse kernel), reached from a different codebase.

What I cannot claim: the fleet has spent the last three days in a benchmark campaign (dozens of short boots, no long uptime), and the watchdog that would have logged a wedge was not loaded during it. So "no wedge observed on B12X" is true and nearly worthless as evidence. The next real number is wedge-free hours under sustained mixed traffic on B12X with the capture armed, and I will report it here either way, including the RAS delta on the first wedge if one comes.

On credit: noted, and agreed. If we port anything from the jasl deepseek_v4 Triton lineage the credit travels with it in-file, the same way we are carrying 0xdfi's overlays with per-file provenance.