Skip to content

[CUDA] QMoE: select the 64-row CTA tile for repacked INT3 prefill when experts average >= 24 rows #33714

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
Tianlei Wu (tianleiwu) merged 1 commit into from
Oct 29, 2026
Merged

[CUDA] QMoE: select the 64-row CTA tile for repacked INT3 prefill when experts average >= 24 rows #33714

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
10 changes: 8 additions & 2 deletions onnxruntime/contrib_ops/cuda/llm/moe_gemm/moe_int2_prefill.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -314,9 +314,15 @@ void RunPackedGroupedGemm(const Int2MoePrefillParams& params,
ORT_ENFORCE(weight_bits == 4, "Packed INT prefill requires INT2 or INT4 weights");
GroupedGemmInput<ElementType, cutlass::uint4b_t, ElementType, ElementType> gemm{
activations, offsets + 1, reinterpret_cast<const cutlass::uint4b_t*>(weights), static_cast<const ElementType*>(scales), nullptr, biases, output, nullptr, nullptr, ActivationType::Identity, num_rows, num_columns, reduction_size, params.num_experts, params.block_size, true, false, params.stream, {}, {}};
// Same crossover as PickPrefillTactic in moe_quantization.cc: the 64-row CTA wins once experts average >= 24 rows.
const double rows_per_expert =
params.num_experts > 0 ? static_cast<double>(num_rows) / static_cast<double>(params.num_experts) : 0.0;
const auto tile_config =
rows_per_expert >= 24.0
? cutlass_extensions::CutlassTileConfig::CtaShape64x128x64_WarpShape64x32x64
: cutlass_extensions::CutlassTileConfig::CtaShape32x128x64_WarpShape32x32x64;
gemm.gemm_config = cutlass_extensions::CutlassGemmConfig(
cutlass_extensions::CutlassTileConfig::CtaShape32x128x64_WarpShape32x32x64,
cutlass_extensions::SplitKStyle::NO_SPLIT_K, 1, 4);
tile_config, cutlass_extensions::SplitKStyle::NO_SPLIT_K, 1, 4);
MoeGemmRunner<ElementType, cutlass::uint4b_t, ElementType> runner(params.sm, params.multiprocessor_count);
runner.moeGemm(gemm, {});
}
Expand Down
4 changes: 4 additions & 0 deletions onnxruntime/test/contrib_ops/moe_test.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3433,6 +3433,10 @@ void RunInt3AndMixedWidthCudaPrefillParityTest() {
// Partial final row tile (16 + 3 = 19 rows with row_tile_size = 16)
{3, 3, 64, 19, 16},
{3, 4, 32, 19, 16},
// 64 rows * top_k 2 / 4 experts = 32 rows per expert: the packed INT4 grouped GEMM picks the 64-row CTA tile
{3, 3, 64, 64, 0},
{3, 4, 32, 64, 0},
{4, 3, 128, 64, 0},
};

for (const auto& tc : cases) {
Expand Down
Loading