Skip to content

Commit 99928b9

Browse files
ElizaWszolarasmith
authored andcommitted
[Bugfix][Kernel] Fix moe align block issue for mixtral (vllm-project#12413)
1 parent 326e481 commit 99928b9

File tree

1 file changed

+3
-1
lines changed

1 file changed

+3
-1
lines changed

csrc/moe/moe_align_sum_kernels.cu

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -33,7 +33,9 @@ __global__ void moe_align_block_size_kernel(scalar_t* __restrict__ topk_ids,
3333

3434
extern __shared__ int32_t shared_mem[];
3535
int32_t* cumsum = shared_mem; // 1d tensor with shape (num_experts + 1)
36-
token_cnts_t* tokens_cnts = (token_cnts_t*)(shared_mem + blockDim.x + 1);
36+
token_cnts_t* tokens_cnts =
37+
(token_cnts_t*)(shared_mem + num_experts +
38+
1); // 2d tensor with shape (blockDim.x + 1, num_experts)
3739

3840
for (int i = 0; i < num_experts; ++i) {
3941
tokens_cnts[index(num_experts, threadIdx.x + 1, i)] = 0;

0 commit comments

Comments
 (0)