vLLM #51593：负长度触发 Persistent Top-K CTA 死锁
DeepSeek-V4 MTP在SM120下的persistent_topk死锁分析
| 项目 | 内容 |
|---|---|
| 上游 Issue | vllm-project/vllm#51593 |
| 触发条件 | DeepSeek-V4、MTP、FULL CUDA Graph padding、SM120 |
| 表面症状 | shm_broadcast timeout、EngineDeadError、GPU 利用率接近 100% |
| 实际故障点 | persistent_topk 的 multi-CTA radix barrier |
根因：
MTP + FULL CUDA Graph padding产生负的per-token context length。persistent_topk将int32 -1作为uint32_t读取后得到UINT_MAX，导致同一CTA group的leader和peer对是否进入multi-CTA radix path作出相反判断。peer提前退出，leader永久等待inter-CTA barrier，最终表现为shm_broadcast timeout和EngineDeadError
复现配置
vllm serve <path-to>/DeepSeek-V4-Flash \
--tensor-parallel-size 4 \
--attention-backend FLASHINFER_MLA_SPARSE_DSV4 \
--moe-backend flashinfer_cutlass \
--speculative-config '{"method":"mtp","num_speculative_tokens":1}' \
--max-num-seqs 16 --max-num-batched-tokens 8192 \
--max-model-len 65536 --kv-cache-dtype fp8 --block-size 256 \
--gpu-memory-utilization 0.9 --tokenizer-mode deepseek_v4
故障链路总览
CUDA Graph padding slot
-> C4 indexer produces context length -1
-> persistent_topk converts -1 to UINT_MAX
-> CTA1 exits based on the host max_seq_len
-> CTA0 enters the radix path based on UINT_MAX
-> CTA0 waits forever for two CTAs at the barrier
-> the GPU kernel never retires
-> the output-copy CUDA event never completes
-> no worker response is enqueued
-> EngineCore eventually reports shm_broadcast timeout
- 实际生成的 padding context lengths: [-1, 0]
+ 必须满足的合法值: [ 0, 0]
根因分析
1. MTP 与 CUDA Graph padding
MTP 是整条故障链的起点。普通 decode 一次 forward 中，每个 request 只处理当前位置的一个 query token，因此 next_n=1。这个复现使用 num_speculative_tokens=1，所以 target verification 一次会处理两个连续位置，也就是 next_n=2：第一个是当前 target position，第二个是 MTP speculative position
-
服务刚开始时可以有很多并发
request，但每个request的输出长度不同，有些会先结束，因此running batch会随着生成过程逐渐下降。所谓剩下 3 个真实request，并不是系统固定创建 3 个request，而只是某一轮scheduler调度时，前面的请求已经陆续结束，此时还有 3 个真实request继续生成 -
当剩下 3 个
request时，MTP一轮实际上需要处理 3 × 2 = 6 个真实query token。但是FULL CUDA Graph不能像eager模式一样每一步都使用任意shape，它会replay已经提前capture好的固定shape。这个场景使用的是8-token graph，所以 6 个真实token必须补到 8 个。额外的两个token正好对应一个完整的MTP request slot，因此graph metadata中实际上会出现 4 个request slot：前三个是真实request，最后一个只是paddingText 3 real requests × 2 MTP positions = 6 real query tokens ↓ FULL CUDA Graph shape = 8 ↓ append 2 padding tokens ↓ 4 request slots query_start_loc = [0, 2, 4, 6, 6] seq_lens = [L0, L1, L2, 0] decode_lens = [2, 2, 2, 0]
2. 合法的 padding 元数据如何变成 -1
-
query_start_loc表示每个request在query-token buffer中的边界，所以前三个request分别对应[0,2)、[2,4)、[4,6)，最后一个padding request对应[6,6)，说明它实际上没有任何真实query token。seq_lens是request级别的当前sequence length，padding request没有历史有效序列，所以它的seq_len=0是完全正确的；decode_lens=0也同样正确。因此CUDA Graph padding本身没有制造非法数据，真正的bug是后面的MTP metadata preparation把这个合法的 0 继续展开成负数 -
native MTP path需要把request-levelseq_lens转成每个speculative position对应的context length。对于一个正常长度为L的request，一次处理两个连续位置时，第一个位置能够看到L-1个KV，第二个位置能够看到L个KV，因此代码使用seq_lens.unsqueeze(1) - max_decode_len + 1 + offsets来生成[L-1, L]。对于正常request这个计算没有问题 -
padding request的seq_len=0仍然被按照固定max_decode_len=2展开，于是变成0 - 2 + 1 + [0,1] = [-1,0]。也就是说，这里第一次产生真正非法的数据：context length语义上表示当前 query 可以访问多少 KV token，它不可能小于 0，padding request正确结果应该是[0,0]Python if use_native and next_n > 1: assert self.decode_seq_lens_buffer.dim() == 1 seq_lens_buffer = self.decode_seq_lens_buffer[ : num_decodes * max_decode_len ].view(num_decodes, max_decode_len) seq_lens_buffer[:] = ( seq_lens.unsqueeze(1) - max_decode_len + 1 + self.offsets_buffer[:max_decode_len] ) seq_lens = seq_lens_buffer return seq_lens, block_table, decode_lens, num_decodes, requires_padding -
普通
decode的next_n=1，同一个公式退化成seq_len - 1 + 1 = seq_len，所以padding的 0 仍然是 0；MTP的next_n=2才会让第一个位置相当于seq_len-1，因此当padding request的seq_len=0时得到-1 -
variable-length flatten path又是另外一种情况：它根据真实decode_lens去展开token，padding request的decode_len=0，因此根本不会生成任何expanded token，剩余buffer还会被清零，所以那条路径不会产生这个-1 -
真正有问题的是
native MTP的固定二维metadata layout，以及公式本身同样未clamp的uniform speculative path
3. 负长度穿过 C4 indexer
DeepSeek-V4的C4 indexer还会把这些context length转换到压缩后的KV空间，compress_ratio=4。正常长度例如 100 会变成 25，但Python的//是floor division，所以-1 // 4仍然是-1，而不是 0。因此这个非法值经过C4后没有被消掉，继续向下流入top-k。与此同时，服务配置里的max_model_len=65536在C4 indexer中实际只对应 65536 / 4 = 16384 个candidate，所以后面persistent_topk看到的logits row width是 16384
4. Leader 与 peer 的判断发生分裂
-
SM120上的sparse indexer不会走cooperative_topk，这一架构被cooperative path明确排除，最终调用的是persistent_topk。这里真正危险的不是单纯收到一个-1，而是kernel内部有两套决定是否进入multi-CTA radix的判断来源。non-leader CTA在kernel很早的位置只检查host已经传进来的params.max_seq_len；当前值是 16384，小于RADIX_THRESHOLD=32768，因此non-leader CTA会认为所有row都是short/medium row，没有必要参与multi-CTA radixText FULL CUDA Graph 需要固定 8-token shape ↓ 3 个真实 MTP requests 只有 6 个 token ↓ 补出第 4 个 padding request ↓ request-level metadata: seq_len = 0 decode_len = 0 ↓ native MTP 按固定 next_n=2 构造 per-token context lengths ↓ 0 → [-1,0] ↓ DeepSeek-V4 C4 compression [-1,0] // 4 → [-1,0] ↓ 非法 -1 进入 persistent_topk -
leader CTA后面处理具体row时却不是看这个host scalar，而是重新从device memory读取params.lengths[row_idx]。问题是lengths的元素本来是int32，但代码直接写成const uint32_t seq_len = params.lengths[row_idx] -
于是
padding row中的-1在任何判断之前就变成UINT_MAX，即 4294967295 -
leader接下来比较seq_len <= RADIX_THRESHOLD时自然得到false，因此它认为这个row是一个超长row，必须进入multi-CTAradix_topk。这时同一个CTA group已经发生逻辑分裂：peer CTA根据max_seq_len=16384提前退出，leader CTA根据错误的row length=UINT_MAX进入只有完整CTA group才能运行的radix pathText padding request ↓ request-level seq_len = 0 ← 合法 ↓ native MTP per-token expansion ↓ [-1, 0] ← 第一次产生非法值 ↓ C4 compression ↓ [-1, 0] ↓ persistent_topk non-leader CTA leader CTA │ │ params.max_seq_len = 16384 lengths[row] = -1 │ │ 16384 <= 32768 int32 → uint32 │ │ early return UINT_MAX │ > 32768 │ radix_topk │ waits for peer CTA forever
5. Inter-CTA barrier 永久等待
radix_topk 进入后使用 arrival_counter 做 inter-CTA barrier。当前 ctas_per_group=2，所以 barrier 初始阶段需要两个 CTA 都到达，target 是 2
但现场 cuda-gdb 抓到的是 arrival_counter=1、target_val=2：只有 leader CTA 到达过，另一个 CTA 已经执行 early return。因为没有任何 CTA 再能把 counter 从 1 加到 2，所以 leader 永远 spin 在 wait_ge()，整个 kernel 永远不会 retire。GPU 因此显示接近 100% utilization，但几乎没有 memory activity；CPU 侧真正接触 GPU 结果的是 async output-copy thread，它一直卡在 copy_event.synchronize()，所以 worker main thread 看起来只是正常 idle 在 zmq_poll
调试现场的决定性证据是
arrival_counter=1、target_val=2：leader CTA已到达barrier，而已经early return的peer CTA永远不会再把计数器推进到 2
修复建议
建议同时修复两层：
producer保证不产生负长度，consumer则维护0 <= seq_len <= min(stride, max_seq_len)，避免其他caller再次破坏kernel invariant
Producer：禁止产生负 context length：修复首先应该从 producer 做，因为 context length 本身就不应该允许负数。native MTP path 应该在构造 per-token sequence lengths 时直接 clamp 到 0；uniform speculative path 使用同样的算式，也应该同时做 lower-bound clamp。这样 padding request 的 [−1,0] 会直接变成 [0,0]，对所有正常 request 的 [L-1,L] 完全没有影响
seq_lens_buffer[:] = (
seq_lens.unsqueeze(1)
- max_decode_len
+ 1
+ self.offsets_buffer[:max_decode_len]
).clamp_min_(0)
Consumer：加固 persistent_topk：第二层应该 harden persistent_topk。kernel 不应该假设所有 caller 永远传入合法 length，尤其这个 length 会直接决定内存访问范围和 cooperative topology。不能先把它读进 uint32_t 再 clamp，因为 -1 此时已经变成 UINT_MAX；应该先保持 signed 类型判断下界，再把合法正数限制到实际 row width
const int32_t raw_seq_len = params.lengths[row_idx];
const uint32_t max_valid_len =
min(params.stride, params.max_seq_len);
const uint32_t seq_len =
raw_seq_len <= 0
? 0u
: min(static_cast<uint32_t>(raw_seq_len),
max_valid_len);
kernel invariant：每个row都满足 0 ≤seq_len≤min(stride, max_seq_len)。这样一旦host侧确认max_seq_len <= RADIX_THRESHOLD，任何具体row也不可能突然大于RADIX_THRESHOLD，从结构上消除peer CTA early return、leader CTA 进入 radix这一整类deadlock。同时upper bound还能防止未来其他caller传入超过logits row width的错误length后产生潜在OOB read
背景：CTA 协作机制
在 CUDA 里，CTA（Cooperative Thread Array）基本可以直接理解成一个 CUDA thread block。一个 kernel launch 会启动很多 block，每个 block 内有很多线程；这个 persistent_topk 里每个 CTA 有 1024 个线程。普通情况下，一个 CTA 可以独立处理一个 row，但当 row 很长时，一个 CTA 的计算和 shared memory 不够高效，所以 persistent_topk 会把多个 CTA 组成一个 CTA group，共同处理同一个 large row。当前这个现场 ctas_per_group=2，意思就是一个 group 里有两个 block：一个是 cta_in_group=0，可以叫 leader；另一个是 cta_in_group=1，可以叫 peer。两者在 large-row radix 路径中不是各算各的，而是会共同构造 histogram、共享全局状态，并在几个阶段通过 arrival_counter 做 barrier 同步。因此只要 leader 进入 multi-CTA radix，peer 就必须也活着并进入同一条路径；否则 barrier 一定等不到
-
persistent_topk为省资源做一个优化：如果整个batch根本没有large row，就没必要让peer CTA一直存在。 因此kernel刚开始时会先看一个全局上界params.max_seq_lenC++ / CUDA const uint32_t ctas_per_group = params.ctas_per_group; const uint32_t cta_in_group = blockIdx.x % ctas_per_group; if (cta_in_group != 0 && params.max_seq_len <= RADIX_THRESHOLD) { return; } -
params.max_seq_len是launch时host已经知道的这个 logits tensor 的最大有效 row width / batch-level upper bound。当前DeepSeek-V4 C4 indexer的logits width是 16384，而RADIX_THRESHOLD=32768，所以kernel得出一个很合理的结论：这个batch里理论上不可能有长度超过 32768 的row。于是group里的peer CTA，也就是cta_in_group=1，直接退出 -
leader CTA留下来，因为short/medium row仍然需要它处理 -
每个实际
row的长度都必须不大于params.max_seq_len。如果全局已经知道最大长度只有 16384，那么后面不应该突然出现一个row长度大于 32768。只要这个条件成立，让peer提前退出就是完全安全的，因为leader后面也绝不可能进入需要peer协作的radix路径。但是leader继续遍历具体row时，又会进行第二次判断C++ / CUDA const uint32_t seq_len = params.lengths[row_idx]; if (seq_len <= RADIX_THRESHOLD) { // 只有 CTA0 处理 short / medium row ... continue; } // 否则认为是 large row radix_topk(...); -
padding MTP row的实际lengths[row_idx]是int32的-1，但代码却直接用uint32_tC++ / CUDA const uint32_t seq_len = params.lengths[row_idx]; -
因此
-1并没有保持为非法负长度，而是按照二进制补码被解释成4294967295。于是同一个CTA group对同一个batch得到两个完全矛盾的结论-
第一次判断看到
params.max_seq_len=16384，所以peer CTA认为没有 large row，可以退出 -
第二次判断中
leader看到padding row的seq_len=4294967295，所以认为这是一个超大 row，必须进入 multi-CTA radix
-
-
进入
radix_topk后，两个CTA原本必须在第一个barrier汇合。每个CTA的thread 0会把arrival_counter加 1，然后等待它达到ctas_per_groupC++ / CUDA if (tx == 0) { red_release(&state->arrival_counter, 1); } wait_ge( &state->arrival_counter, (barrier_phase + 1) * ctas_per_group, tx ); -
当前
ctas_per_group=2，所以正常应该是leader加 1、peer再加 1，counter到 2 后两边继续执行。但peer已经在第一个判断那里return，因此现在只有leader能把counter加到 1。cuda-gdb现场看到的正好就是arrival_counter=1、target_val=2
persistent_topk 原本允许 peer CTA 根据全局 max_seq_len 提前退出，因为它假设具体 row length 永远不会超过这个全局上界；MTP padding 产生的 -1 被错误转成 UINT_MAX 后，这个假设被破坏，导致 peer 根据第一个判断退出，而 leader 根据第二个判断进入必须依赖 peer 的 cooperative radix，最终形成不可恢复的 barrier deadlock
背景：Radix Top-K
在 vLLM 中，persistent_topk 主要用于对每个 query 对应的一行候选 logits 进行 Top-K 选择，即从大量 KV candidate 中筛选出分数最高的 K 个位置。它并不要求对整行元素进行完整排序，而只需要确定哪些 candidate 属于前 K 名，因此本质上属于 selection problem，而不是 full sorting problem。这类操作常见于稀疏注意力、候选路由或大规模 KV 筛选场景中，可以显著减少后续需要参与计算的候选数量
- 对于较长的
logits row，vLLM采用Radix Top-K / Radix Select。其核心思想是将FP32 score转换成保持数值顺序的32-bit integer key，然后按高位到低位逐步定位第K大元素。32-bit key被划分为 4 个8-bit段，每一轮根据当前8 bit将候选划分到 256 个bucket中，并通过从高bucket到低bucket的累计计数，判断第K大元素落在哪个bucket。确定该bucket后，只保留具有相同高位前缀的候选，在下一轮继续检查更低的8 bit。经过 4 轮后，就能够得到第K大元素对应的完整32-bit key，即Top-K的边界值pivot
struct PersistentTopKParams {
const float* input; // [num_rows, stride]
int32_t* output; // [num_rows, top_k]
const int32_t* lengths; // [num_rows]
RadixRowState* row_states;
uint32_t num_rows;
uint32_t stride;
uint32_t top_k;
uint32_t chunk_size;
uint32_t ctas_per_group;
uint32_t max_seq_len;
};
struct RadixRowState {
uint32_t histogram[3][256];
uint32_t remaining_k;
uint32_t prefix;
int arrival_counter;
int output_counter;
};
因此，这个算法的关键并不是生成一个完全有序的 Top-K 序列，而是高效地确定 Top-K 分界点。所有分数严格大于 pivot 的元素必然属于 Top-K，而与 pivot 相等的元素只需要选取足够数量以补满 K 个结果。相比对长度为 N 的 logits 完整排序，Radix Select 只需要进行少量固定轮次的线性扫描和 histogram 统计，更适合 GPU 上处理数万甚至更长的候选序列
- 在
vLLM的实现中，超长row还会被划分给多个CTA并行处理。每个CTA负责一段logits，独立统计局部256-bin histogram，再将结果合并成整行的全局histogram；随后所有CTA基于相同的统计结果共同确定下一轮的radix prefix。因而，这套实现可以概括为：将长logits row分片并行扫描，通过 4 轮8-bit histogram逐步确定第K大元素的数值边界，再根据该边界输出最终Top-K candidate indices。 它的主要价值在于避免全排序，并充分利用GPU的并行histogram和多CTA协同能力来降低大规模Top-K选择的成本
CTA shared memory
┌────────────────────────────┐
│ local_histogram[256] │ 当前 CTA 自己 chunk 的 histogram
├────────────────────────────┤
│ suffix_sum[256] │ 合并后的 histogram 做 suffix scan
├────────────────────────────┤
│ shared_scalars[5] │
│ [0] prefix │
│ [1] remaining_k │
│ [2] selected_bucket │
│ [3] next_remaining_k │
│ [...] │
├────────────────────────────┤
│ shared_ordered[chunk_size] │ 本 CTA 负责的 FP32 ordered keys
└────────────────────────────┘