Skip to content

[vLLM][MoE] Update kernel grid to use autotune selected BLOCK_M and BLOCK_N values #7776

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

Open
sarah12121212 wants to merge 20 commits into
base: main
Choose a base branch
from
+6 −16
Open

[vLLM][MoE] Update kernel grid to use autotune selected BLOCK_M and BLOCK_N values #7776

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
20 commits
Select commit Hold shift + click to select a range
2b82f56
update
sarah12121212 Aug 14, 2026
cc65619
Merge branch 'main' into moe2-batched-moe2
sarah12121212 Aug 17, 2026
b1024b4
Update vllm-tests-reusable.yml
sarah12121212 Aug 17, 2026
2cb2fc3
Fix batched MoE TD autotune cache key
sarah12121212 Aug 17, 2026
d9ff27e
Fix batched MoE autotuned launch grid
sarah12121212 Aug 18, 2026
2529e66
Remove unrelated SiluAndMul reference changes
sarah12121212 Aug 18, 2026
ea372b7
Revert "Fix batched MoE TD autotune cache key"
sarah12121212 Aug 18, 2026
2cad8af
Merge branch 'main' into batched-moe-td-autotune-key
sarah12121212 Aug 18, 2026
45386c0
Merge branch 'main' into batched-moe-td-autotune-key
sarah12121212 Aug 18, 2026
9dbde47
Merge branch 'main' into batched-moe-td-autotune-key
sarah12121212 Aug 19, 2026
733e9dc
Merge remote-tracking branch 'origin/main' into batched-moe-td-autotu…
sarah12121212 Aug 20, 2026
f6ea1b3
Merge branch 'main' into batched-moe-td-autotune-key
sarah12121212 Aug 20, 2026
a0836c9
Merge branch 'main' into batched-moe-td-autotune-key
sarah12121212 Aug 21, 2026
0ae5223
Update vllm-tests-reusable.yml
sarah12121212 Aug 21, 2026
95fb40e
Update vllm_tdesc.txt
sarah12121212 Aug 21, 2026
aef6c3b
Update vllm_tdesc.txt
sarah12121212 Aug 21, 2026
55f791f
Update vllm_moe.txt
sarah12121212 Aug 24, 2026
57ad399
Update vllm_moe.txt
sarah12121212 Aug 24, 2026
af956bb
Merge branch 'main' into batched-moe-td-autotune-key
sarah12121212 Aug 24, 2026
42a8b3b
update leftover patch noise
sarah12121212 Aug 24, 2026
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
9 changes: 6 additions & 3 deletions benchmarks/triton_kernels_benchmark/vllm/batched_moe/batched_moe.patch
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -165,12 +165,15 @@ index 609b28a95..6ceb8d9ff 100644
num_pid_n = tl.cdiv(N, BLOCK_N)
pid_m = pid_mn // num_pid_n
pid_n = pid_mn % num_pid_n
@@ -449,8 +487,8 @@ def invoke_moe_batched_triton_kernel(
@@ -449,8 +487,9 @@ def invoke_moe_batched_triton_kernel(
BLOCK_K = config["BLOCK_SIZE_K"]

grid = (
- grid = (
- expert_num_tokens.size(0),
triton.cdiv(max_num_tokens, BLOCK_M) * triton.cdiv(B.size(1), BLOCK_N),
- triton.cdiv(max_num_tokens, BLOCK_M) * triton.cdiv(B.size(1), BLOCK_N),
+ grid = lambda meta: (
+ triton.cdiv(max_num_tokens, meta["BLOCK_M"])
+ * triton.cdiv(B.size(1), meta["BLOCK_N"]),
Comment thread
quinnlp marked this conversation as resolved.
+ expert_num_tokens.size(0),
)

Expand Down
7 changes: 0 additions & 7 deletions scripts/skiplist/default/vllm_tdesc.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -882,10 +882,3 @@ tests/kernels/attention/test_triton_unified_attention.py::test_triton_unified_at
tests/kernels/attention/test_triton_unified_attention.py::test_triton_unified_attn[8-q_dtype1-2048-50.0-dtype0-256-16-256-num_heads1-seq_lens1]
tests/kernels/attention/test_triton_unified_attention.py::test_triton_unified_attn[8-q_dtype1-2048-50.0-dtype0-256-16-256-num_heads2-seq_lens0]
tests/kernels/attention/test_triton_unified_attention.py::test_triton_unified_attn[8-q_dtype1-2048-50.0-dtype0-256-16-256-num_heads2-seq_lens1]

# Tensor-likes are not close: https://github.com/intel/intel-xpu-backend-for-triton/issues/7753
tests/kernels/moe/test_batched_moe.py::test_fused_moe_batched_experts[False-None-False-dtype0-2-8-222-128-2048]
tests/kernels/moe/test_batched_moe.py::test_fused_moe_batched_experts[False-None-False-dtype0-6-8-45-1024-128]
tests/kernels/moe/test_batched_moe.py::test_fused_moe_batched_experts[False-None-False-dtype0-6-8-64-512-512]
tests/kernels/moe/test_batched_moe.py::test_fused_moe_batched_experts[False-None-False-dtype0-6-8-64-1024-2048]
tests/kernels/moe/test_batched_moe.py::test_fused_moe_batched_experts[False-None-False-dtype0-6-8-222-128-2048]
6 changes: 0 additions & 6 deletions scripts/skiplist/xe2/vllm_tdesc.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -885,9 +885,3 @@ tests/kernels/attention/test_triton_unified_attention.py::test_triton_unified_at

# Nondeterministic tensor-likes are not close on B580: https://github.com/intel/intel-xpu-backend-for-triton/issues/7395
tests/kernels/attention/test_triton_unified_attention.py::test_triton_unified_attn_fp16_input_fp8_output[8-32768-50.0-None-16-128-num_heads1-seq_lens3]

# Tensor-likes are not close: https://github.com/intel/intel-xpu-backend-for-triton/issues/7753
tests/kernels/moe/test_batched_moe.py::test_fused_moe_batched_experts[False-None-False-dtype0-2-8-222-128-2048]
tests/kernels/moe/test_batched_moe.py::test_fused_moe_batched_experts[False-None-False-dtype0-6-8-45-1024-128]
tests/kernels/moe/test_batched_moe.py::test_fused_moe_batched_experts[False-None-False-dtype0-6-8-64-512-512]
tests/kernels/moe/test_batched_moe.py::test_fused_moe_batched_experts[False-None-False-dtype0-6-8-222-128-2048]
Loading