EngineX-Hygon/sglang
5
0
Fork 0
Code Issues Pull Requests Actions 7 Projects Releases Wiki Activity

[BugFix] fix pre_reorder_triton_kernel default int32 issue (#7814)

Browse Source
This commit is contained in:
Morpheus Guo
2025-07-13 04:42:36 +08:00
committed by GitHub
parent 98aa836bbf
commit 5f6756b038
1 changed files with 4 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

6
python/sglang/srt/layers/moe/ep_moe/kernels.py
View File

@@ -236,7 +236,8 @@ def pre_reorder_triton_kernel(
): ):
OutDtype = gateup_input_ptr.dtype.element_ty OutDtype = gateup_input_ptr.dtype.element_ty
src_idx = tl.program_id(0) src_idx_int32 = tl.program_id(0)
src_idx = src_idx_int32.to(tl.int64)
src2dst_ptr = src2dst_ptr + src_idx * topk src2dst_ptr = src2dst_ptr + src_idx * topk
topk_ids_ptr = topk_ids_ptr + src_idx * topk topk_ids_ptr = topk_ids_ptr + src_idx * topk
src_ptr = input_ptr + src_idx * hidden_size src_ptr = input_ptr + src_idx * hidden_size
@@ -255,7 +256,8 @@ def pre_reorder_triton_kernel(
else: else:
scale = 1.0 scale = 1.0
dst_idx = tl.load(src2dst_ptr + idx) dst_idx_int32 = tl.load(src2dst_ptr + idx)
dst_idx = dst_idx_int32.to(tl.int64)
dst_ptr = gateup_input_ptr + dst_idx * hidden_size dst_ptr = gateup_input_ptr + dst_idx * hidden_size
for start_offset in tl.range(0, hidden_size, BLOCK_SIZE): for start_offset in tl.range(0, hidden_size, BLOCK_SIZE):
offset = start_offset + vec offset = start_offset + vec