From 2e53f9725b4aed93909eada64f0066cb7cc7275f Mon Sep 17 00:00:00 2001
From: coderfeli <coderfeli@163.com>
Date: Mon, 10 Feb 2025 01:26:08 +0000
Subject: [PATCH] skip empty expert

---
 .../gridwise_gemm_xdl_cshuffle_v3_multi_d_b_preshuffle.hpp   | 5 ++++-
 1 file changed, 4 insertions(+), 1 deletion(-)

diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d_b_preshuffle.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d_b_preshuffle.hpp
index b78433bb30..b51231c0eb 100644
--- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d_b_preshuffle.hpp
+++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d_b_preshuffle.hpp
@@ -1126,7 +1126,10 @@ struct GridwiseGemmMultiD_xdl_cshuffle_v3_b_preshuffle
         const index_t m_block_data_idx_on_grid =
             __builtin_amdgcn_readfirstlane(block_m_id * MPerBlock);
         const index_t expert_stride = __builtin_amdgcn_readfirstlane(problem.N * problem.K);
-
+                
+        const index_t t0 = (p_sorted_token_ids[block_m_id * MPerBlock] & 0xffffff);
+        if(t0 >= problem.NumTokens)
+            return;
         // N0, K0, Blocksize*KPack
         const index_t n_block_data_idx_on_grid =
             __builtin_amdgcn_readfirstlane(block_n_id * NXdlPerWave);