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 32128ff53d..8f8fc205ec 100644 --- a/benchmarks/triton_kernels_benchmark/vllm/batched_moe/batched_moe.patch +++ b/benchmarks/triton_kernels_benchmark/vllm/batched_moe/batched_moe.patch @@ -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"]), + expert_num_tokens.size(0), ) 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] 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]