Skip to content

Commit

Permalink
[Bugfix] Make moe_align_block_size AMD-compatible (vllm-project#3470)
Browse files Browse the repository at this point in the history
  • Loading branch information
WoosukKwon authored Mar 18, 2024
1 parent e8efbe8 commit 097b8b5
Showing 1 changed file with 2 additions and 1 deletion.
3 changes: 2 additions & 1 deletion csrc/moe_align_block_size_kernels.cu
Original file line number Diff line number Diff line change
Expand Up @@ -111,7 +111,8 @@ void moe_align_block_size(

// set dynamic shared mem
auto kernel = vllm::moe_align_block_size_kernel<scalar_t>;
AT_CUDA_CHECK(cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, shared_mem));
AT_CUDA_CHECK(
VLLM_DevFuncAttribute_SET_MaxDynamicSharedMemorySize((void *)kernel, shared_mem));
kernel<<<1, num_experts, shared_mem, stream>>>(
topk_ids.data_ptr<scalar_t>(),
sorted_token_ids.data_ptr<int32_t>(),
Expand Down

0 comments on commit 097b8b5

Please sign in to comment.