Skip to content

Add CUTLASS FP8 Blockscale MoE kernel for Hopper architecture #7278

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 20 commits into from
Jul 3, 2025
Merged

Add CUTLASS FP8 Blockscale MoE kernel for Hopper architecture #7278

Changes from 1 commit
Commits
Show all changes
20 commits
Select commit Hold shift + click to select a range
1af8a4b
feat/support sm90_fp8_blockwise_scaled_group_mm
ayrnb Jun 10, 2025
bbe9198
fix bug
ayrnb Jun 10, 2025
5fadc62
Fix sm_count for H20 GPU.
HydraQYH Jun 11, 2025
4ea113b
fix/code refine
ayrnb Jun 11, 2025
aeeca14
[benchmark]: single kernel perf of cutlass/deepgemm/triton group gemm
TianQiLin666666 Jun 12, 2025
f9e55be
Use ATen to get device & properties.
HydraQYH Jun 12, 2025
465e8f5
Fix benchmark.
HydraQYH Jun 13, 2025
7587bba
Optimize MMAConfig2 by using Pingpong schedule and new TiledMMA shape…
HydraQYH Jun 14, 2025
488ff0f
[benchmark] add more shapes
TianQiLin666666 Jun 16, 2025
117ecb8
add ep cases
TianQiLin666666 Jun 17, 2025
47e7404
Use unique config for sm90 fp8 blockwise scaling grouped gemm.
HydraQYH Jun 17, 2025
54af9e9
Refine code. Support unitest.
HydraQYH Jun 17, 2025
33fb8b9
add Qwen3 benchmark cases
TianQiLin666666 Jun 18, 2025
8b80ac7
fix bench shapes
TianQiLin666666 Jun 18, 2025
3d81974
Remove duplicated check.
HydraQYH Jun 18, 2025
5cfea52
Merge branch 'main' into feat/sm90_fp8_blockwise_grouped_gemm_r1
FlamingoPg Jul 1, 2025
28f6545
update init
ayrnb Jul 1, 2025
8d737dd
pre-commit
ayrnb Jul 1, 2025
559fc54
format
ayrnb Jul 1, 2025
95f2860
Merge branch 'main' into feat/sm90_fp8_blockwise_grouped_gemm_r1
ayrnb Jul 1, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
7 changes: 5 additions & 2 deletions sgl-kernel/csrc/moe/fp8_blockwise_moe_kernel.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -120,8 +120,10 @@ void launch_sm90_fp8_blockwise_scaled_group_mm(
reinterpret_cast<typename ScheduleConfig::LayoutSFB*>(layout_sfb.data_ptr())};

cutlass::KernelHardwareInfo hw_info;
// TODO(qiyuhang): get device_id by cudaGetDevice
hw_info.device_id = 0;
hw_info.sm_count = 132;
// TODO(qiyuhang): get sm_count by cudaGetDeviceProperties
hw_info.sm_count = 78; // H20 config
Copy link
Collaborator

@yuan-luo yuan-luo Jul 23, 2025

Choose a reason for hiding this comment

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

This kernel sm_count configuration is not good in H800.

$python3 sgl-kernel/benchmark/bench_fp8_blockwise_group_gemm.py

Benchmark: expected_m_per_group=128, n=512, k=7168, num_groups=256
deepgemm: 493.7056064605713 us
cutlass: 460.3487968444824 us

Benchmark: expected_m_per_group=256, n=512, k=7168, num_groups=256
deepgemm: 563.2063865661621 us
cutlass: 671.1967945098877 us

Benchmark: expected_m_per_group=256, n=256, k=7168, num_groups=256
deepgemm: 391.4144039154053 us
cutlass: 397.8463888168335 us

Benchmark: expected_m_per_group=512, n=256, k=7168, num_groups=256
deepgemm: 598.5184192657471 us
cutlass: 743.3504104614258 us

Copy link
Collaborator

@HydraQYH HydraQYH Jul 23, 2025

Choose a reason for hiding this comment

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

@yuan-luo What is the SM Count on H800? It seems that this may be caused by load imbalance. We have some engineers working on optimizing load balancing. Since I lack the H800 machine, could you please provide me with an ncu report?

Copy link
Collaborator

@HydraQYH HydraQYH Jul 23, 2025

Choose a reason for hiding this comment

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

@yuan-luo We use SM Count as the number of CTAs. I checked and found that H800 has 144 SM Cores, the same as H100. In the second and fourth configurations, there are only (256/128)*(512/128)*256 = 2048 Output Tiles. 2048 / 144 = 14.22 The last stage may only enable 20% of the SM Cores. I guess this may be the cause of the performance issue and hope you can help provide an ncu report to confirm.

Copy link
Collaborator

@yuan-luo yuan-luo Aug 4, 2025

Choose a reason for hiding this comment

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

@HydraQYH I'm working on a related PR about FP8 MoE kernel for Hopper. Will update and relate with this PR later on.


typename GemmKernel::EpilogueArguments epilogue_args{
{},
Expand All @@ -140,7 +142,8 @@ void launch_sm90_fp8_blockwise_scaled_group_mm(

at::cuda::CUDAGuard device_guard{(char)a_ptrs.get_device()};
const cudaStream_t stream = at::cuda::getCurrentCUDAStream(a_ptrs.get_device());


// TODO(qiyuhang): skip can_implement when problem_sizes_host is nullptr
auto can_implement_status = gemm_op.can_implement(args);
TORCH_CHECK(can_implement_status == cutlass::Status::kSuccess, "Failed to implement GEMM");

Expand Down