Skip to content

[Feature] Support fusedmoe on Blackwell - #5325

Merged
zoooo0820 merged 8 commits into
PaddlePaddle:developfrom
Echo-Nie:sm100
Dec 16, 2025
Merged

zoooo0820 merged 8 commits into
PaddlePaddle:developfrom
Echo-Nie:sm100

Conversation

@Echo-Nie

@Echo-Nie Echo-Nie commented Dec 2, 2025

Copy link
Copy Markdown
Contributor

Motivation

Support fusedmoe on Blackwell

Modifications

prerequisite PR: Paddle#76708
custom_ops/gpu_ops/cutlass_kernels/moe_gemm/fused_moe_cutlass_kernel.h
custom_ops/gpu_ops/cutlass_kernels/moe_gemm/fused_moe_gemm_kernels_template.h

Usage or Command

In test_fusedmoe.py, add the following import:

from fastdeploy.model_executor.layers.quantization.weight_only import WeightOnlyConfig, WINT4Config, WINT8Config

Then uncomment the corresponding lines:

quant_config = BlockWiseFP8Config(weight_block_size=[128, 128])
# quant_config = WINT8Config({})
# quant_config = WINT4Config({})

Finally, run:

python FastDeploy/tests/layers/test_fusedmoe.py

Accuracy Tests

BF16

num_token: 4096
[13.9 13.9 13.9 13.9 13.9]
13.9 57.9
0.6 TB/s
num_token: 8192
[27.2 27.2 27.2 27.2 27.2]
27.2 59.2
0.3 TB/s
num_token: 16384
[53.8 53.8 53.8 53.8 53.8]
53.8 59.9
0.1 TB/s
num_token: 32768
[106.9 106.9 106.9 106.9 106.9]
106.9 60.3
0.1 TB/s

BlockWiseFP8Config

num_token: 4096
[9.7 9.7 9.7 9.7 9.7]
9.7 83.0
0.8 TB/s
num_token: 8192
[18.1 18.1 18.1 18.1 18.1]
18.1 89.0
0.4 TB/s
num_token: 16384
[34.8 34.7 34.8 34.7 34.8]
34.8 92.6
0.2 TB/s
num_token: 32768
[68.  68.1 68.1 68.1 68.1]
68.1 94.6
0.1 TB/s

WINT8Config

num_token: 4096
[14.4 14.4 14.4 14.4 14.4]
14.4 55.9
0.6 TB/s
num_token: 8192
[28.3 28.3 28.3 28.3 28.3]
28.3 56.9
0.3 TB/s
num_token: 16384
[55.8 55.9 55.8 55.9 55.8]
55.8 57.7
0.1 TB/s
num_token: 32768
[111.1 111.1 111.1 111.1 111.1]
111.1 58.0
0.1 TB/s

WINT4Config

num_token: 4096
[14.8 14.8 14.8 14.8 14.8]
14.8 54.4
0.5 TB/s
num_token: 8192
[28.9 28.9 28.9 28.9 28.9]
28.9 55.7
0.3 TB/s
num_token: 16384
[57.3 57.3 57.3 57.3 57.2]
57.2 56.3
0.1 TB/s
num_token: 32768
[113.9 113.9 113.9 113.9 113.8]
113.8 56.6
0.1 TB/s

Checklist

  • Add at least a tag in the PR title.
    • Tag list: [[FDConfig],[APIServer],[Engine], [Scheduler], [PD Disaggregation], [Executor], [Graph Optimization], [Speculative Decoding], [RL], [Models], [Quantization], [Loader], [OP], [KVCache], [DataProcessor], [BugFix], [Docs], [CI], [Optimization], [Feature], [Benchmark], [Others], [XPU], [HPU], [GCU], [DCU], [Iluvatar], [Metax]]
    • You can add new tags based on the PR content, but the semantics must be clear.
  • Format your code, run pre-commit before commit.
  • Add unit tests. Please write the reason in this PR if no unit tests.
  • Provide accuracy results.
  • If the current PR is submitting to the release branch, make sure the PR has been submitted to the develop branch, then cherry-pick it to the release branch with the [Cherry-Pick] PR tag.

@codecov-commenter

codecov-commenter commented Dec 2, 2025

Copy link
Copy Markdown

Codecov Report

✅ All modified and coverable lines are covered by tests.
⚠️ Please upload report for BASE (develop@532f9ba). Learn more about missing BASE report.

Additional details and impacted files
@@            Coverage Diff             @@
##             develop    #5325   +/-   ##
==========================================
  Coverage           ?   60.87%           
==========================================
  Files              ?      329           
  Lines              ?    41160           
  Branches           ?     6274           
==========================================
  Hits               ?    25055           
  Misses             ?    14213           
  Partials           ?     1892           
Flag Coverage Δ
GPU 60.87% <ø> (?)

Flags with carried forward coverage won't be shown. Click here to find out more.

☔ View full report in Codecov by Sentry.
📢 Have feedback on the report? Share it here.

🚀 New features to boost your workflow:
  • ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.

@paddle-bot

paddle-bot Bot commented Dec 2, 2025

Copy link
Copy Markdown

Thanks for your contribution!

KernelRunner<WintQuantMethod::kWeightOnlyInt2, true>::run_kernel(params, shared_storage);
#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 800) && (__CUDA_ARCH__ < 1010)
KernelRunner<WintQuantMethod::kWeightOnlyInt2, true>::run_kernel(
params, shared_storage);

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

910 改成 1010,放行 SM100

} else if (sm_ >= 75 && sm_ < 80) {
dispatch_moe_gemm_to_cutlass_macro(cutlass::arch::Sm75);
} else if (sm_ >= 80 && sm_ < 91) {
} else if (sm_ >= 80 && sm_ < 101) {

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

91 改成 101,放行 SM100

@zoooo0820 zoooo0820 left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM

@zoooo0820
zoooo0820 merged commit 50100f9 into PaddlePaddle:develop Dec 16, 2025
15 of 18 checks passed
@Echo-Nie
Echo-Nie deleted the sm100 branch December 16, 2025 05:24
chang-wenbin pushed a commit to chang-wenbin/FastDeploy that referenced this pull request Mar 2, 2026
xiaoguoguo626807 pushed a commit to xiaoguoguo626807/FastDeploy that referenced this pull request May 7, 2026
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