Skip to content
GitLab
Menu
Projects
Groups
Snippets
Loading...
Help
Help
Support
Community forum
Keyboard shortcuts
?
Submit feedback
Contribute to GitLab
Sign in
Toggle navigation
Menu
Open sidebar
OpenDAS
vllm_cscc
Commits
221d388c
Unverified
Commit
221d388c
authored
Jan 24, 2025
by
ElizaWszola
Committed by
GitHub
Jan 25, 2025
Browse files
[Bugfix][Kernel] Fix moe align block issue for mixtral (#12413)
parent
3132a933
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
3 additions
and
1 deletion
+3
-1
csrc/moe/moe_align_sum_kernels.cu
csrc/moe/moe_align_sum_kernels.cu
+3
-1
No files found.
csrc/moe/moe_align_sum_kernels.cu
View file @
221d388c
...
@@ -33,7 +33,9 @@ __global__ void moe_align_block_size_kernel(scalar_t* __restrict__ topk_ids,
...
@@ -33,7 +33,9 @@ __global__ void moe_align_block_size_kernel(scalar_t* __restrict__ topk_ids,
extern
__shared__
int32_t
shared_mem
[];
extern
__shared__
int32_t
shared_mem
[];
int32_t
*
cumsum
=
shared_mem
;
// 1d tensor with shape (num_experts + 1)
int32_t
*
cumsum
=
shared_mem
;
// 1d tensor with shape (num_experts + 1)
token_cnts_t
*
tokens_cnts
=
(
token_cnts_t
*
)(
shared_mem
+
blockDim
.
x
+
1
);
token_cnts_t
*
tokens_cnts
=
(
token_cnts_t
*
)(
shared_mem
+
num_experts
+
1
);
// 2d tensor with shape (blockDim.x + 1, num_experts)
for
(
int
i
=
0
;
i
<
num_experts
;
++
i
)
{
for
(
int
i
=
0
;
i
<
num_experts
;
++
i
)
{
tokens_cnts
[
index
(
num_experts
,
threadIdx
.
x
+
1
,
i
)]
=
0
;
tokens_cnts
[
index
(
num_experts
,
threadIdx
.
x
+
1
,
i
)]
=
0
;
...
...
Write
Preview
Markdown
is supported
0%
Try again
or
attach a new file
.
Attach a file
Cancel
You are about to add
0
people
to the discussion. Proceed with caution.
Finish editing this message first!
Cancel
Please
register
or
sign in
to comment