From 2b82f563be8e72c67909c59a5938693e5a240bdd Mon Sep 17 00:00:00 2001 From: Sarah Date: Fri, 14 Aug 2026 20:26:13 +0000 Subject: [PATCH 01/12] update --- scripts/skiplist/default/vllm_moe.txt | 1 - scripts/skiplist/xe2/vllm_moe.txt | 1 - scripts/vllm/vllm-fix.patch | 49 ++++++++++++++++++++++++++- 3 files changed, 48 insertions(+), 3 deletions(-) diff --git a/scripts/skiplist/default/vllm_moe.txt b/scripts/skiplist/default/vllm_moe.txt index c9962fe549..ffb950a3bc 100644 --- a/scripts/skiplist/default/vllm_moe.txt +++ b/scripts/skiplist/default/vllm_moe.txt @@ -5,7 +5,6 @@ tests/kernels/moe/test_moe.py::test_fused_moe_int64_overflow # Waiting for upstream fix https://github.com/intel/intel-xpu-backend-for-triton/issues/7705 tests/kernels/moe/test_silu_mul_fp8_quant_deep_gemm.py::test_silu_mul_fp8_quant_deep_gemm -tests/kernels/moe/test_batched_moe.py::test_fused_moe_batched_experts tests/kernels/moe/test_moe.py::test_fused_moe_wn16 # Accuracy error on PVC: https://github.com/intel/intel-xpu-backend-for-triton/issues/7131 diff --git a/scripts/skiplist/xe2/vllm_moe.txt b/scripts/skiplist/xe2/vllm_moe.txt index 493553ccc2..0bdbfa8405 100644 --- a/scripts/skiplist/xe2/vllm_moe.txt +++ b/scripts/skiplist/xe2/vllm_moe.txt @@ -5,7 +5,6 @@ tests/kernels/moe/test_moe.py::test_fused_moe_int64_overflow # Waiting for upstream fix https://github.com/intel/intel-xpu-backend-for-triton/issues/7705 tests/kernels/moe/test_silu_mul_fp8_quant_deep_gemm.py::test_silu_mul_fp8_quant_deep_gemm -tests/kernels/moe/test_batched_moe.py::test_fused_moe_batched_experts tests/kernels/moe/test_moe.py::test_fused_moe_wn16 # Block FP8 modular MoE tests require FP32-SiLU reference https://github.com/intel/intel-xpu-backend-for-triton/issues/7614 diff --git a/scripts/vllm/vllm-fix.patch b/scripts/vllm/vllm-fix.patch index ea6b75e7d8..3a53cc943c 100644 --- a/scripts/vllm/vllm-fix.patch +++ b/scripts/vllm/vllm-fix.patch @@ -98,7 +98,7 @@ index b82c4d82c3..483672927a 100644 out += (x * D).to(out.dtype) out = (out if z is None else out * F.silu(z)).to(x.dtype) diff --git a/tests/kernels/moe/test_moe.py b/tests/kernels/moe/test_moe.py -index a519980dea..a51a91de5c 100644 +index a519980dea..51db9360aa 100644 --- a/tests/kernels/moe/test_moe.py +++ b/tests/kernels/moe/test_moe.py @@ -334,6 +334,20 @@ def test_fused_moe( @@ -170,6 +170,53 @@ index a519980dea..a51a91de5c 100644 def test_batched_fused_marlin_moe( m: int, n: int, +diff --git a/tests/kernels/utils.py b/tests/kernels/utils.py +index cc1d1bbf8..7fc766d8a 100644 +--- a/tests/kernels/utils.py ++++ b/tests/kernels/utils.py +@@ -914,6 +914,8 @@ def torch_experts( + f32 = torch.float32 + + act = op_registry[activation.custom_op_name] ++ silu_act_fn = SiluAndMul(compile_native=False) ++ act_fn = silu_act_fn if act is SiluAndMul else act() + + for i in range(num_experts): + mask = topk_ids == i +@@ -922,13 +924,13 @@ def torch_experts( + tmp1 = a[mask] @ w1[i].transpose(0, 1) + if b_bias1 is not None: + tmp1 = tmp1 + b_bias1[i].view(1, -1).to(tmp1.dtype) +- tmp2 = act()(tmp1) ++ tmp2 = act_fn(tmp1) + out[mask] = tmp2 @ w2[i].transpose(0, 1) + if b_bias2 is not None: + out[mask] = out[mask] + b_bias2[i].view(1, -1).to(tmp1.dtype) + elif quant_input_only: + tmp1 = a[mask] @ w1[i].transpose(0, 1) +- tmp2 = SiluAndMul()(tmp1) ++ tmp2 = silu_act_fn(tmp1) + tmp2, tmp2_scale = moe_kernel_quantize_input( + tmp2, None, quant_dtype, per_act_token_quant + ) +@@ -946,7 +948,7 @@ def torch_experts( + ) + if b_bias1 is not None: + tmp1 = tmp1 + b_bias1[i].view(1, -1).to(tmp1.dtype) +- tmp2 = SiluAndMul()(tmp1) ++ tmp2 = silu_act_fn(tmp1) + tmp2, b_scale = moe_kernel_quantize_input( + tmp2, a2_scale, quant_dtype, per_act_token_quant, block_shape + ) +@@ -970,7 +972,7 @@ def torch_experts( + if b_bias1 is not None: + tmp1 = tmp1 + b_bias1[i].view(1, -1).to(out.dtype) + +- tmp2 = act()(tmp1).to(out.dtype) ++ tmp2 = act_fn(tmp1).to(out.dtype) + + tmp2, b_scale = moe_kernel_quantize_input( + tmp2, a2_scale, quant_dtype, per_act_token_quant, block_shape diff --git a/tests/utils.py b/tests/utils.py index 90f12d444e..1be15bd7f4 100644 --- a/tests/utils.py From b1024b4b433ef3799c8ce56bf2ecd5473d37bffd Mon Sep 17 00:00:00 2001 From: Sarah Date: Mon, 17 Aug 2026 13:19:52 -0600 Subject: [PATCH 02/12] Update vllm-tests-reusable.yml --- .github/workflows/vllm-tests-reusable.yml | 8 -------- 1 file changed, 8 deletions(-) diff --git a/.github/workflows/vllm-tests-reusable.yml b/.github/workflows/vllm-tests-reusable.yml index ed9ba10202..fdddda4d0c 100644 --- a/.github/workflows/vllm-tests-reusable.yml +++ b/.github/workflows/vllm-tests-reusable.yml @@ -134,15 +134,7 @@ jobs: fail-fast: false matrix: suite: - - vllm-spec-decode - - vllm-mrv2 - vllm-moe - - vllm-triton-attn - - vllm-mamba - - vllm-quant - - vllm-kda - - vllm-tdesc - - vllm-rest runs-on: ${{ fromJSON((inputs.runner_label == '' || startsWith(inputs.runner_label, 'max')) && format('["linux", "rolling", "{0}"]', inputs.runner_label || 'max1100') || format('["linux", "{0}"]', inputs.runner_label)) }} timeout-minutes: 180 defaults: From 2cb2fc3bb06d8c31ce45e684c1d558dc52fe9483 Mon Sep 17 00:00:00 2001 From: sarah Date: Mon, 17 Aug 2026 22:58:37 +0000 Subject: [PATCH 03/12] Fix batched MoE TD autotune cache key --- .../triton_kernels_benchmark/vllm/batched_moe/batched_moe.patch | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/benchmarks/triton_kernels_benchmark/vllm/batched_moe/batched_moe.patch b/benchmarks/triton_kernels_benchmark/vllm/batched_moe/batched_moe.patch index 73a3280bc4..5234ce132a 100644 --- a/benchmarks/triton_kernels_benchmark/vllm/batched_moe/batched_moe.patch +++ b/benchmarks/triton_kernels_benchmark/vllm/batched_moe/batched_moe.patch @@ -64,7 +64,7 @@ index 609b28a95..6ceb8d9ff 100644 + +@triton.autotune( + configs=get_batched_moe_configs(), -+ key=['max_num_tokens', 'K', 'N'], ++ key=['max_num_tokens', 'K', 'N', 'USE_TD'], +) @triton.jit def batched_triton_kernel( From d9ff27ef3214d9f0feb6d84270d547332eea2f41 Mon Sep 17 00:00:00 2001 From: sarah Date: Tue, 18 Aug 2026 00:02:24 +0000 Subject: [PATCH 04/12] Fix batched MoE autotuned launch grid --- .../vllm/batched_moe/batched_moe.patch | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/benchmarks/triton_kernels_benchmark/vllm/batched_moe/batched_moe.patch b/benchmarks/triton_kernels_benchmark/vllm/batched_moe/batched_moe.patch index 5234ce132a..e7b6322f16 100644 --- a/benchmarks/triton_kernels_benchmark/vllm/batched_moe/batched_moe.patch +++ b/benchmarks/triton_kernels_benchmark/vllm/batched_moe/batched_moe.patch @@ -135,12 +135,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"]), + expert_num_tokens.size(0), ) From 2529e66a7181466d4d39b40afb38757e90172a17 Mon Sep 17 00:00:00 2001 From: sarah Date: Tue, 18 Aug 2026 14:36:57 +0000 Subject: [PATCH 05/12] Remove unrelated SiluAndMul reference changes --- scripts/vllm/vllm-fix.patch | 47 ------------------------------------- 1 file changed, 47 deletions(-) diff --git a/scripts/vllm/vllm-fix.patch b/scripts/vllm/vllm-fix.patch index 3a53cc943c..206784ea2d 100644 --- a/scripts/vllm/vllm-fix.patch +++ b/scripts/vllm/vllm-fix.patch @@ -170,53 +170,6 @@ index a519980dea..51db9360aa 100644 def test_batched_fused_marlin_moe( m: int, n: int, -diff --git a/tests/kernels/utils.py b/tests/kernels/utils.py -index cc1d1bbf8..7fc766d8a 100644 ---- a/tests/kernels/utils.py -+++ b/tests/kernels/utils.py -@@ -914,6 +914,8 @@ def torch_experts( - f32 = torch.float32 - - act = op_registry[activation.custom_op_name] -+ silu_act_fn = SiluAndMul(compile_native=False) -+ act_fn = silu_act_fn if act is SiluAndMul else act() - - for i in range(num_experts): - mask = topk_ids == i -@@ -922,13 +924,13 @@ def torch_experts( - tmp1 = a[mask] @ w1[i].transpose(0, 1) - if b_bias1 is not None: - tmp1 = tmp1 + b_bias1[i].view(1, -1).to(tmp1.dtype) -- tmp2 = act()(tmp1) -+ tmp2 = act_fn(tmp1) - out[mask] = tmp2 @ w2[i].transpose(0, 1) - if b_bias2 is not None: - out[mask] = out[mask] + b_bias2[i].view(1, -1).to(tmp1.dtype) - elif quant_input_only: - tmp1 = a[mask] @ w1[i].transpose(0, 1) -- tmp2 = SiluAndMul()(tmp1) -+ tmp2 = silu_act_fn(tmp1) - tmp2, tmp2_scale = moe_kernel_quantize_input( - tmp2, None, quant_dtype, per_act_token_quant - ) -@@ -946,7 +948,7 @@ def torch_experts( - ) - if b_bias1 is not None: - tmp1 = tmp1 + b_bias1[i].view(1, -1).to(tmp1.dtype) -- tmp2 = SiluAndMul()(tmp1) -+ tmp2 = silu_act_fn(tmp1) - tmp2, b_scale = moe_kernel_quantize_input( - tmp2, a2_scale, quant_dtype, per_act_token_quant, block_shape - ) -@@ -970,7 +972,7 @@ def torch_experts( - if b_bias1 is not None: - tmp1 = tmp1 + b_bias1[i].view(1, -1).to(out.dtype) - -- tmp2 = act()(tmp1).to(out.dtype) -+ tmp2 = act_fn(tmp1).to(out.dtype) - - tmp2, b_scale = moe_kernel_quantize_input( - tmp2, a2_scale, quant_dtype, per_act_token_quant, block_shape diff --git a/tests/utils.py b/tests/utils.py index 90f12d444e..1be15bd7f4 100644 --- a/tests/utils.py From ea372b72c1aae854a4c3d5d7f15cecb6d4a29a40 Mon Sep 17 00:00:00 2001 From: sarah Date: Tue, 18 Aug 2026 14:43:50 +0000 Subject: [PATCH 06/12] Revert "Fix batched MoE TD autotune cache key" This reverts commit 2cb2fc3bb06d8c31ce45e684c1d558dc52fe9483. --- .../triton_kernels_benchmark/vllm/batched_moe/batched_moe.patch | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/benchmarks/triton_kernels_benchmark/vllm/batched_moe/batched_moe.patch b/benchmarks/triton_kernels_benchmark/vllm/batched_moe/batched_moe.patch index e7b6322f16..29bbb8f480 100644 --- a/benchmarks/triton_kernels_benchmark/vllm/batched_moe/batched_moe.patch +++ b/benchmarks/triton_kernels_benchmark/vllm/batched_moe/batched_moe.patch @@ -64,7 +64,7 @@ index 609b28a95..6ceb8d9ff 100644 + +@triton.autotune( + configs=get_batched_moe_configs(), -+ key=['max_num_tokens', 'K', 'N', 'USE_TD'], ++ key=['max_num_tokens', 'K', 'N'], +) @triton.jit def batched_triton_kernel( From 0ae522336d77b2cdd21f7943c24d7fc3bc50fc71 Mon Sep 17 00:00:00 2001 From: Sarah Date: Fri, 21 Aug 2026 08:33:29 -0600 Subject: [PATCH 07/12] Update vllm-tests-reusable.yml --- .github/workflows/vllm-tests-reusable.yml | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/.github/workflows/vllm-tests-reusable.yml b/.github/workflows/vllm-tests-reusable.yml index fdddda4d0c..ed9ba10202 100644 --- a/.github/workflows/vllm-tests-reusable.yml +++ b/.github/workflows/vllm-tests-reusable.yml @@ -134,7 +134,15 @@ jobs: fail-fast: false matrix: suite: + - vllm-spec-decode + - vllm-mrv2 - vllm-moe + - vllm-triton-attn + - vllm-mamba + - vllm-quant + - vllm-kda + - vllm-tdesc + - vllm-rest runs-on: ${{ fromJSON((inputs.runner_label == '' || startsWith(inputs.runner_label, 'max')) && format('["linux", "rolling", "{0}"]', inputs.runner_label || 'max1100') || format('["linux", "{0}"]', inputs.runner_label)) }} timeout-minutes: 180 defaults: From 95fb40ef1d0d9d4bab6a298f8cbe99d8dcf4f13f Mon Sep 17 00:00:00 2001 From: Sarah Date: Fri, 21 Aug 2026 09:19:23 -0600 Subject: [PATCH 08/12] Update vllm_tdesc.txt --- scripts/skiplist/xe2/vllm_tdesc.txt | 6 ------ 1 file changed, 6 deletions(-) diff --git a/scripts/skiplist/xe2/vllm_tdesc.txt b/scripts/skiplist/xe2/vllm_tdesc.txt index 01bcc19d8b..b7c9648b19 100644 --- a/scripts/skiplist/xe2/vllm_tdesc.txt +++ b/scripts/skiplist/xe2/vllm_tdesc.txt @@ -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] From aef6c3b8233d423a3a85cf71239dfeba9ea127f3 Mon Sep 17 00:00:00 2001 From: Sarah Date: Fri, 21 Aug 2026 09:19:47 -0600 Subject: [PATCH 09/12] Update vllm_tdesc.txt --- scripts/skiplist/default/vllm_tdesc.txt | 7 ------- 1 file changed, 7 deletions(-) diff --git a/scripts/skiplist/default/vllm_tdesc.txt b/scripts/skiplist/default/vllm_tdesc.txt index cab2cfb13b..89252fda68 100644 --- a/scripts/skiplist/default/vllm_tdesc.txt +++ b/scripts/skiplist/default/vllm_tdesc.txt @@ -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] From 55f791fc6814273faf499052cfa14796c3ed0d8b Mon Sep 17 00:00:00 2001 From: Sarah Date: Mon, 24 Aug 2026 06:40:13 -0600 Subject: [PATCH 10/12] Update vllm_moe.txt --- scripts/skiplist/default/vllm_moe.txt | 1 + 1 file changed, 1 insertion(+) diff --git a/scripts/skiplist/default/vllm_moe.txt b/scripts/skiplist/default/vllm_moe.txt index 5a70cc2198..2dfc366aed 100644 --- a/scripts/skiplist/default/vllm_moe.txt +++ b/scripts/skiplist/default/vllm_moe.txt @@ -2,6 +2,7 @@ # Large-token fused MoE accuracy failure on XPU after the memory-query patch. tests/kernels/moe/test_moe.py::test_fused_moe_int64_overflow +tests/kernels/moe/test_batched_moe.py::test_fused_moe_batched_experts tests/kernels/moe/test_moe.py::test_fused_moe_wn16 # Accuracy error on PVC: https://github.com/intel/intel-xpu-backend-for-triton/issues/7131 From 57ad399965b2fa9086e075e9fc3f44d9aca6f7ea Mon Sep 17 00:00:00 2001 From: Sarah Date: Mon, 24 Aug 2026 06:40:59 -0600 Subject: [PATCH 11/12] Update vllm_moe.txt --- scripts/skiplist/xe2/vllm_moe.txt | 1 + 1 file changed, 1 insertion(+) diff --git a/scripts/skiplist/xe2/vllm_moe.txt b/scripts/skiplist/xe2/vllm_moe.txt index 3565710942..ba0a0af197 100644 --- a/scripts/skiplist/xe2/vllm_moe.txt +++ b/scripts/skiplist/xe2/vllm_moe.txt @@ -2,6 +2,7 @@ # Large-token fused MoE accuracy failure on XPU after the memory-query patch. tests/kernels/moe/test_moe.py::test_fused_moe_int64_overflow +tests/kernels/moe/test_batched_moe.py::test_fused_moe_batched_experts tests/kernels/moe/test_moe.py::test_fused_moe_wn16 # Block FP8 modular MoE tests require FP32-SiLU reference https://github.com/intel/intel-xpu-backend-for-triton/issues/7614 From 42a8b3b29f476ac18f9047f65e867b631736b221 Mon Sep 17 00:00:00 2001 From: sarah Date: Mon, 24 Aug 2026 13:04:20 +0000 Subject: [PATCH 12/12] update leftover patch noise --- scripts/vllm/vllm-fix.patch | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/scripts/vllm/vllm-fix.patch b/scripts/vllm/vllm-fix.patch index 206784ea2d..ea6b75e7d8 100644 --- a/scripts/vllm/vllm-fix.patch +++ b/scripts/vllm/vllm-fix.patch @@ -98,7 +98,7 @@ index b82c4d82c3..483672927a 100644 out += (x * D).to(out.dtype) out = (out if z is None else out * F.silu(z)).to(x.dtype) diff --git a/tests/kernels/moe/test_moe.py b/tests/kernels/moe/test_moe.py -index a519980dea..51db9360aa 100644 +index a519980dea..a51a91de5c 100644 --- a/tests/kernels/moe/test_moe.py +++ b/tests/kernels/moe/test_moe.py @@ -334,6 +334,20 @@ def test_fused_moe(