Skip to content

Feat/adapt sm120 rtx5000pro nvfp4 - #1223

Merged
LLLLKKKK merged 7 commits into
alibaba:feat/adapt-sm120-rtx5000profrom
ykykzq:feat/adapt-sm120-rtx5000pro
Aug 7, 2026
Merged

LLLLKKKK merged 7 commits into
alibaba:feat/adapt-sm120-rtx5000profrom
ykykzq:feat/adapt-sm120-rtx5000pro

Conversation

@ykykzq

@ykykzq ykykzq commented Jul 24, 2026

Copy link
Copy Markdown

feat(sm120): 接入 flashinfer b12x,为 sm_120/121 添加 NVFP4 MoE 算子支持

背景

在RTX Pro 5000(SM 12.0)上以 NVFP4精度运行 MoE 模型(如 Qwen3.5-35B-A3B-NVFP4)时,rtp-llm 现有的两条 NVFP4 MoE 路径在 sm_120/121 上暂时均不可用:

  • TrtllmFp4Executor(trtllm-gen cubin):trtllm-gen 预编译 cubin 仅覆盖 SM100 系列架构(SM100/103/110),SM120 上没有对应 MoE cubin。
  • CutedslFp4Executor(CuTe DSL blockscaled,tcgen05):kernel 同样依赖 SM100 独有的 tcgen05 指令。

flashinfer 已为 sm_120/121 专门提供b12x CuTe DSL fused MoE,但是其目前限制CUDA版本须>=13。经本地实验验证,该限制过于严格,实际上CUDA 12.9 已经可以对sm_120f架构进行正确编译。其执行、数值全部正确。本 PR 据此把 b12x 接入 rtp-llm。

改动内容

  1. 新增 B12xFp4Executor.../executors/b12x_fp4_executor.py),在其中添加ptach放宽cuda版本要求
  2. 策略选型和自动 fallback(strategy/fp4.pystrategy/__init__.pyfused_moe/__init__.pydefs/type.py
  3. 权重加载分支(device_impl.py
  4. 启动参数(server_args/moe_group_args.py

测试

新增 .../executors/test/b12x_fp4_executor_test.py,对算子进行数值正确性单测:

  • 按 b12x loader 完全一致的方式构造 NVFP4 权重,将 fused 输出与 bf16 dequant 参考 MoE 对比。
  • 以余弦相似度判定(global cos ≥ 0.98、worst-row cos ≥ 0.95)算子精度正确性,实测global cos 0.986, worst-row cos 0.984。

相关内容

  • flashinfer #3066(b12x CuTe DSL nvfp4 fused MoE)、#3670(trtllm-gen FP4在 SM120上 不支持)。

Loading
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants