[Bug] Qwen4Exp QSA indexer: per-chunk logits buffer grows with max_seq_len, caching allocator keeps every size, device OOM/hang on unified-m

问题场景

用户在 2× NVIDIA DGX Spark（GB10、SM121、unified memory）上以 TP=2 多节点方式运行 vLLM，加载 Qwen4Exp 架构模型 `RadixArk/Qwen3.8-Flash-Next-NVFP4`，启用 chunked prefill、prefix caching、MTP k=3、`cudagraph_mode FULL_DECODE_ONLY`，对约 254K token 的长 prompt 做冷 prefill 时触发故障。故障表现为步骤在正好 166,400 computed tokens 处挂起，驱动日志出现 OOM，约 300 秒后 EngineCore 因 RPC 超时死亡。该问题涉及 `vllm/models/qwen4_exp/nvidia/ops/qsa_indexer.py` 中的 `qsa_select_paged_prefill` 代码路径（上游未修改部分）。

报错原文

NVRM: nvCheckOkFailedNoLog: Check failed: Out of memory [NV_ERR_NO_MEMORY] (0x00000051) returned from _memdescAllocInternal(pMemDesc) @ mem_desc.c:1359

TimeoutError: RPC call to sample_tokens timed out.

SchedulerOutput: num_computed_tokens=[..., 166400], num_scheduled_tokens={...: 3200}

NotImplementedError: Qwen4Exp QSA requires a BF16 main KV cache

说明：最后一条 `NotImplementedError` 出现于复现者按正文参数启动时，属另一独立问题——正文命令中的 `–kv-cache-dtype fp8` 在该版本上会被 `Qwen4ExpQSA.__init__` 拒绝，因此不能用于复现本故障。

原因分析

最可能的原因在 `qsa_select_paged_prefill` 中：该函数现在按 batch 的 `max_seq_len` 计算 indexer logits buffer 宽度：

logits_width = triton.cdiv(triton.cdiv(max_seq_len, compress_ratio), 64) * 64
logits_width = min(max(64, logits_width), page_table.shape[1] * k_cache.shape[1])
max_logits_bytes = envs.VLLM_SPARSE_INDEXER_MAX_LOGITS_MB * 1024 * 1024   # 512 MB default
rows_per_chunk = max(1, max_logits_bytes // (logits_width * 4))
...
logits = torch.empty((num_queries, logits_width), dtype=torch.float32, ...)

而在 commit `a5c9179e73`（PR #54915，2026-09-04）之前，宽度为固定的 page-table 宽度（`columns = page_table.shape[1] * k_cache.shape[1]`），每块 chunk 都分配同一 512 MB buffer，caching allocator 可复用同一块。

改为随 `max_seq_len` 增长后，在 chunked prefill（每 chunk 3200 token）配合 `compress_ratio=4` 时，第 i 个 chunk 的 buffer 约为 `3200 × (800·i) × 4 B = 10.24 MB × i`，随 chunk 单调递增；当 `i = 52` 时 `52 × 3200 = 166,400` token，正好与故障挂起点一致。由于每块尺寸不同，PyTorch caching allocator 找不到足够大的已缓存块，每个 chunk 都触发新的 `cudaMalloc` 且保留旧块，累计约 `Σ 10.24 MB × i (i ≤ 52) ≈ 14 GB per rank`。

在独立 GPU 上，`cudaMalloc` 最终会失败，allocator 会清理缓存并重试；但在 unified-memory 的 GB10 上，`cudaMalloc` 会一直成功直到驱动无内存可用，故障以 GPU hang（dmesg 中 `NV_ERR_NO_MEMORY`）而非 Python OOM 的形式出现。

不过需要标注：单 GB10 复现测试（含 250,010 token prefill，默认 512 MB 预算）全部完成，因此纯 buffer 增长可能不足以独立解释该故障。复现者指出 TP=2 下权重被拆分，每节点可用内存应更多而非更少，因此问题也可能与多节点/TP 路径，或实际启动时的内存余量有关。

环境排查

• 确认 vLLM 版本/build 是否落在引入该宽度变更的 commit 之后，特别是 `a5c9179e73`（“Compact indexer logits workspace to improve prefill efficiency”，PR #54915）；报告者确认可用的 nightly 为 `8a728663`（`0.28.1rc1.dev388`），出问题的为 `2a02f6efe`。
• 确认实际启动命令，尤其是 `–kv-cache-dtype` 的取值；在 `2a02f6efe` 上 `vllm/models/qwen4_exp/nvidia/qsa.py:109` 会拒绝 `fp8`，只接受 `auto`/`bfloat16`。
• 记录启动日志中的 `GPU KV cache size:` 与 `Available KV cache memory:` 行，以判断实际内存余量——Issue 中缺失该信息。
• 确认是否使用 TP=2 / 多节点（`–nnodes 2`），以及各节点的可用宿主内存与 unified memory 占用（node-exporter 显示各 rank 在 prefill 最后约 60 秒额外占用 ~9–14 GB，worker 13 → 4 GB 可用，head 9 → 0 GB）。
• 确认模型配置中的 `indexer_compress_ratio=4`、`indexer_budget=2048`，以及 `–max-model-len 262144`、`–max-num-batched-tokens 4096`、chunked prefill、`gpu_memory_utilization 0.8` 等参数。
• 确认是否加载了本地 overlay 或非上游改动，例如 fp8 KV for QSA（#55557）、refusal-direction hook、GDN kernel 修改、PLE-offload connector；报告者称 `qsa_indexer.py` 路径本身为上游未修改。
• 确认硬件与驱动：GB10、SM121、aarch64、CUDA 13、driver 版本，以及是否为 unified memory 机型。

解决步骤

1. 先核对实际启动命令与 vLLM build，确认 `–kv-cache-dtype` 不是 `fp8`（在 `2a02f6efe` 上会直接抛 `NotImplementedError`）。使用实际运行的命令与对应版本的日志进行排障。
2. 在启动日志中记录 `GPU KV cache size:` 与 `Available KV cache memory:`，并核对各 rank 的可用内存余量；这是判断是否因内存余量不足而放大的关键数据。
3. 复现时用 node-exporter 或类似工具监控 prefill 过程中各 rank 在 KV pool 之外的额外内存占用，确认是否随 chunk 单调增长（报告者观察到 prefill 最后约 60 秒额外占用 ~9–14 GB）。
4. 可优先尝试减小 QSA indexer logits workspace 相关的内存预算（例如 `VLLM_SPARSE_INDEXER_MAX_LOGITS_MB`），观察是否延后或避免挂起点；注意复现者已在单 GB10 上以 64 MB 与 512 MB 两种预算跑过 170K/250K prefill，均完成，因此该调整未必能在多节点/TP 场景下解决问题。
5. 关注并验证 PR #56500；报告者称该 PR 在其单 GB10 上可应用、运行于 sm_121 且 170K 无回归，但因其无法复现故障，不能视为对修复的确认。
6. 如果希望在启动前拦截长 prefill 的内存 fit 问题，可参考 Issue 评论中提到的 Badgr 命令（由 @michaelmanly 提出），在启动前检查内存并捕获长 prefill 的 OOM/hang，而不是让工作负载空转。
7. 若环境允许，尝试在单节点/TP=1 上复现，以区分是 buffer 增长本身还是多节点/TP 路径触发的问题。

验证方法

在应用任何规避措施或更新到包含 PR #56500 的版本后，用与报告相同规模的长 prompt（约 254K token、chunked prefill、3200-token chunk）进行冷 prefill，观察是否仍会在 166,400 computed tokens 处挂起，同时确认：

• dmesg 中不再出现 `NVRM: nvCheckOkFailedNoLog: Check failed: Out of memory [NV_ERR_NO_MEMORY] … _memdescAllocInternal(pMemDesc) @ mem_desc.c:1359`；
• EngineCore 不再因 `TimeoutError: RPC call to sample_tokens timed out.` 死亡；
• node-exporter 中 KV pool 之外的额外内存不再随 chunk 单调增长；
• prefill 能完整完成到目标 token 数（例如 250,010 token），且各 rank 内存余量保持稳定。

注意：单 GB10 上报告者的 250K prefill 在默认 512 MB 预算下也能完成（104.8–105.2 s，250,010 tok），因此单机通过并不能证明多节点/TP 场景已修复；请在触发故障的原环境（TP=2 多节点、GB10 unified memory）上验证。

参考来源

vllm-project/vllm #56457

AI 工具推荐

想把多个 AI 模型放在一个入口？

GamsGo AI 集成 ChatGPT、DeepSeek、Gemini、Claude、Midjourney、Veo 等常用模型，适合写作、绘图、视频和日常 AI 工作流。

了解 GamsGo AI

推广链接：通过此链接购买，我可能获得佣金，不影响你的价格。