RuntimeError: Worker failed with error ‘Assertion error

快速结论：该报错通常发生在 NVIDIA GB10（sm_121 架构）设备上尝试启动 DeepSeek-V4 模型时，vLLM 检测到 DeepGEMM 可用但没有正确的架构判断，导致在编译内核中触发断言错误。优先排查是否通过环境变量或补丁正确限制了 DeepGEMM 在 sm_121 上的启用。

适用环境：vLLM 0.26.1rc1.dev608+（aarch64 容器镜像）；2× NVIDIA GB10（sm_121，每节点 121 GiB 统一内存）；CUDA 13.0；驱动 580.173.02；两节点 TP=2 多节点 mp 执行器；DeepSeek-V4-Flash-0731-NVFP4 模型（混合精度，路由专家 NVFP4）。

最快修复方案：合并并应用 PR #53055 提供的修复（为 mhc_pre_broadcast_tilelang 增加 DeepGEMM 架构守卫并回退到 _torch_hc_prenorm_gemm；同时更新 CutlassScaledMM 的 compute_capability 检查）。该 PR 已在双节点 sm_121 环境验证可走到注意力阶段，但后续仍有两个阻塞点需要额外补丁。

注意事项：PR #53055 的验证过程修正了原始诊断——CUTLASS c3x 内核在 sm_121 上实际可用，真正的失败原因是 E8M0 权重缩放格式在 CUDA 上缺少消费者（需将 E8M0 无损上转为 fp32）；此外 deep_gemm_fp8_o_proj 路径仍未加守卫。该 PR 的修复并不完整，需结合后续两个补丁才能完全启动模型。

问题场景

用户在两节点 NVIDIA GB10（sm_121）环境运行 vLLM，通过内置多节点 mp 执行器（–nnodes 2 –distributed-executor-backend mp）启动 DeepSeek-V4-Flash-0731-NVFP4 模型时失败。即使设置了 VLLM_USE_DEEP_GEMM=0 和 VLLM_MOE_USE_DEEP_GEMM=0，模型仍然无法启动。同环境下其他模型（Qwen3-235B-A22B-NVFP4、Mistral-Small-4-119B、Qwen3.8-27B-FP8）均可正常启动，排除集群链路问题。

报错原文

RuntimeError: Worker failed with error 'Assertion error
(/workspace/.deps/deepgemm-src/csrc/apis/hyperconnection.hpp:56): Unsupported architecture'

后续还可能出现：

RuntimeError: dispatch_scaled_mm,
/workspace/csrc/libtorch_stable/quantization/w8a8/cutlass/c3x/scaled_mm_helper.hpp:17

RuntimeError: Assertion error (deepgemm-src/.../utils/layout.hpp:39): t.dim() == N

原因分析

存在多个独立问题叠加导致启动失败：

核心问题 1（原始报告）：mhc_pre_broadcast_tilelang 路径中直接调用 tf32_hc_prenorm_gemm 的第三个调用点未经过 is_deep_gemm_supported() 守卫。该路径不受 VLLM_USE_DEEP_GEMM=0 控制，在 sm_121 上进入 DeepGEMM 内核后触发 Unsupported architecture 断言。
核心问题 2（原始报告的诊断修正）：用户最初认为 CutlassScaledMM.is_supported() 忽略 compute_capability 导致选择错误内核，但证据表明 c3x 内核在 sm_121 上实际运行正常（pertensor 和 blockwise FP8 都能正确输出）。真正原因是该检查点使用 E8M0（ue8m0）权重缩放格式，CUDA 上除 DeepGEMM 外没有其他内核直接消费 E8M0 张量，fp8.py 中只判断 is_scale_e8m0 而未判断 DeepGEMM 是否真的被选中，导致 scaled_mm_helper.hpp:17 的 dtype 检查失败。
阻塞点 4：deep_gemm_fp8_o_proj 路径同样没有架构守卫和无回退方案，在 attention 输出投影阶段触发 layout.hpp:39 断言（t.dim() == N）。
补充确认：DeepGEMM 平台门控问题本身已由 PR #53055 修复（support_deep_gemm() 排除 sm_121 后，权重后处理不再报 Unknown SF transformation 错误）。

环境排查

确认 vLLM 版本是否为 0.26.1rc1.dev608+ 及对应容器镜像；PR #53055 的验证版本为 7b89dfd。
确认 GPU 架构为 sm_121（NVIDIA GB10），可用 nvidia-smi 或 torch.cuda.get_device_capability() 验证。
确认 CUDA 版本 13.0 和驱动 580.173.02 匹配。
确认模型权重格式：quantization_config.scale_fmt 是否为 “ue8m0″（E8M0 格式）。
检查 DeepGEMM 扩展只包含 sm90/sm100 内核：strings vllm/third_party/deep_gemm/_C.cpython-312-aarch64-linux-gnu.so | grep -oE ‘sm[0-9]{2,3}’ | sort -u
确认 VLLM_USE_DEEP_GEMM 和 VLLM_MOE_USE_DEEP_GEMM 设置是否符合预期（注意该环境变量只影响部分调用点）。
检查 VLLM_DISABLED_KERNELS 是否已设置，以及是否能影响 BlockScaledMM 后端选择（该方式不足以解决问题，因为 Triton 同样不直接支持 E8M0）。

解决步骤

应用 PR #53055 的架构守卫修复：为 mhc_pre_broadcast_tilelang 增加 is_deep_gemm_supported() 检查，在 sm_121 上回退到 _torch_hc_prenorm_gemm（该函数已存在且实现相同契约）；同时更新 CutlassScaledMM.is_supported() 使执行在不支持的架构上干净回退到 Triton。合并后模型可加载 48 个分片并越过该点。
处理 E8M0 权重缩放格式问题：在为 Fp8BlockScaledMMLinearKernel 加载权重后，用已有的 _upcast_e8m0_to_fp32 将 E8M0 无损上转为 fp32（E8M0 仅指数位，转换无损失）。该操作在 fp8_utils.py:960 的 Triton 路径在 ROCm/XPU 上已有先例（fp8_utils.py:917），但 CUDA 上需要增加到 process_weights_after_loading。上转换后 CUTLASS 可正常服务该层，无需禁用内核。
为 deep_gemm_fp8_o_proj 添加架构守卫和回退：该路径无保护直接调用 fp8_einsum，在 sm_121 上需绕过或提供等效实现，否则会在注意力输出投影阶段再次中断。
如果在修复完整前需要临时绕过：可设置 VLLM_DISABLED_KERNELS=CutlassFp8BlockScaledMMKernel 让选择回退到 Triton，但需注意 Triton 内核在 CUDA 上同样不支持 E8M0 张量直接绑定，因此该方案不能独立解决问题。

验证方法

在相同的双节点 sm_121 环境、以默认配置（不设置 VLLM_USE_DEEP_GEMM=0）运行目标模型。确认：

profile_run 能顺利通过 mHC broadcast 路径到达注意力块。
权重后处理不再触发 layout.hpp:60 Unknown SF transformation 错误，且 48 个分片全部加载完成（参考基线：Model loading took 77.86 GiB memory and 724.526497 seconds）。
注意力输入投影的 scaled_mm 调用不再报 scaled_mm_helper.hpp:17 的 dtype 错误。
注意力输出投影的 o_proj 不再报 layout.hpp:39 断言错误。
最终模型能正常 start 并 serve，与同环境下其他 NVFP4 模型行为一致。