Skip to content

Make router_fusion to adapt for the large num_of_expert(>2048) #2582

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
Autumn1998 wants to merge 2 commits into
base: main
Choose a base branch
from
+9 −1
Open

Make router_fusion to adapt for the large num_of_expert(>2048) #2582

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
8b77d1c
fix router
Autumn1998 Jan 9, 2026
273a247
[pre-commit.ci] auto fixes from pre-commit.com hooks
pre-commit-ci[bot] Jan 9, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion tests/pytorch/test_fused_router.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -402,7 +402,7 @@ def profile_topk_softmax(
test_topk_softmax(
dtype=torch.float32,
num_tokens=1024,
num_experts=128,
num_experts=3000,
topk=4,
use_pre_softmax=False,
group_topk=None,
Expand Down
4 changes: 4 additions & 0 deletions transformer_engine/common/fused_router/fused_score_for_moe_aux_loss.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -147,6 +147,8 @@ void fused_score_for_moe_aux_loss_forward_kernel_launcher(
size_t shared_memory_size = num_experts * num_token_per_block * sizeof(DataType) // logits
+ topk * num_token_per_block * sizeof(DataType) // topk_logits
+ topk * num_token_per_block * sizeof(int); // topk_indices
cudaFuncSetAttribute(fused_score_for_moe_aux_loss_forward_kernel<DataType>,
cudaFuncAttributeMaxDynamicSharedMemorySize, shared_memory_size);
Comment on lines +150 to +151
Copy link
Contributor

@greptile-apps greptile-apps bot Jan 9, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Following the pattern used elsewhere in the codebase (e.g., ln_fwd_cuda_kernel.cu), this call should be conditional:

Suggested change
cudaFuncSetAttribute(fused_score_for_moe_aux_loss_forward_kernel<DataType>,
cudaFuncAttributeMaxDynamicSharedMemorySize, shared_memory_size);
if (shared_memory_size >= 48 * 1024) {
cudaFuncSetAttribute(fused_score_for_moe_aux_loss_forward_kernel<DataType>,
cudaFuncAttributeMaxDynamicSharedMemorySize, shared_memory_size);
}

This attribute only needs to be set when exceeding the 48KB default limit.

Comment on lines +150 to +151
Copy link
Contributor

@greptile-apps greptile-apps bot Jan 9, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Missing error handling for cudaFuncSetAttribute. The codebase pattern is to wrap CUDA API calls with NVTE_CHECK_CUDA() for proper error handling (see examples in ln_fwd_cuda_kernel.cu, rmsnorm_fwd_cuda_kernel.cu, and other files).

If cudaFuncSetAttribute fails (e.g., if the requested shared memory size exceeds device limits), the error should be caught and reported.

Suggested change
cudaFuncSetAttribute(fused_score_for_moe_aux_loss_forward_kernel<DataType>,
cudaFuncAttributeMaxDynamicSharedMemorySize, shared_memory_size);
NVTE_CHECK_CUDA(cudaFuncSetAttribute(fused_score_for_moe_aux_loss_forward_kernel<DataType>,
cudaFuncAttributeMaxDynamicSharedMemorySize, shared_memory_size));

fused_score_for_moe_aux_loss_forward_kernel<DataType>
<<<grid_size, kThreadsPerBlock, shared_memory_size, stream>>>(
logits, num_tokens, num_experts, topk, score_function, scores, routing_map,
Expand Down Expand Up @@ -283,6 +285,8 @@ void fused_score_for_moe_aux_loss_backward_kernel_launcher(
+
num_experts * num_token_per_block * sizeof(DataType) // act_from_fwd
+ num_experts * num_token_per_block * sizeof(DataType); // comp_buf
cudaFuncSetAttribute(fused_score_for_moe_aux_loss_backward_kernel<DataType>,
cudaFuncAttributeMaxDynamicSharedMemorySize, shared_memory_size);
Comment on lines +288 to +289
Copy link
Contributor

@greptile-apps greptile-apps bot Jan 9, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Following the pattern used elsewhere in the codebase (e.g., ln_fwd_cuda_kernel.cu), this call should be conditional:

Suggested change
cudaFuncSetAttribute(fused_score_for_moe_aux_loss_backward_kernel<DataType>,
cudaFuncAttributeMaxDynamicSharedMemorySize, shared_memory_size);
if (shared_memory_size >= 48 * 1024) {
cudaFuncSetAttribute(fused_score_for_moe_aux_loss_backward_kernel<DataType>,
cudaFuncAttributeMaxDynamicSharedMemorySize, shared_memory_size);
}

This attribute only needs to be set when exceeding the 48KB default limit.

Comment on lines +288 to +289
Copy link
Contributor

@greptile-apps greptile-apps bot Jan 9, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Missing error handling for cudaFuncSetAttribute. The codebase pattern is to wrap CUDA API calls with NVTE_CHECK_CUDA() for proper error handling.

Suggested change
cudaFuncSetAttribute(fused_score_for_moe_aux_loss_backward_kernel<DataType>,
cudaFuncAttributeMaxDynamicSharedMemorySize, shared_memory_size);
NVTE_CHECK_CUDA(cudaFuncSetAttribute(fused_score_for_moe_aux_loss_backward_kernel<DataType>,
cudaFuncAttributeMaxDynamicSharedMemorySize, shared_memory_size));

Note: If this suggestion doesn't match your team's coding style, reply to this and let me know. I'll remember it for next time!

fused_score_for_moe_aux_loss_backward_kernel<DataType>
<<<grid_size, kThreadsPerBlock, shared_memory_size, stream>>>(
intermediate_output, grad_scores, num_tokens, num_experts, topk, score_function,
Expand Down
4 changes: 4 additions & 0 deletions transformer_engine/common/fused_router/fused_topk_with_score_function.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -253,6 +253,8 @@ void fused_topk_with_score_function_forward_kernel_launcher(
shared_memory_size += num_groups * num_token_per_block * sizeof(DataType); // group_scores
shared_memory_size += num_experts * num_token_per_block * sizeof(DataType); // maksed_scores
}
cudaFuncSetAttribute(fused_topk_with_score_function_forward_kernel<DataType, BiasType>,
cudaFuncAttributeMaxDynamicSharedMemorySize, shared_memory_size);
Comment on lines +256 to +257
Copy link
Contributor

@greptile-apps greptile-apps bot Jan 9, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Following the pattern used elsewhere in the codebase (e.g., ln_fwd_cuda_kernel.cu), this call should be conditional:

Suggested change
cudaFuncSetAttribute(fused_topk_with_score_function_forward_kernel<DataType, BiasType>,
cudaFuncAttributeMaxDynamicSharedMemorySize, shared_memory_size);
if (shared_memory_size >= 48 * 1024) {
cudaFuncSetAttribute(fused_topk_with_score_function_forward_kernel<DataType, BiasType>,
cudaFuncAttributeMaxDynamicSharedMemorySize, shared_memory_size);
}

This attribute only needs to be set when exceeding the 48KB default limit.

Comment on lines +256 to +257
Copy link
Contributor

@greptile-apps greptile-apps bot Jan 9, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Missing error handling for cudaFuncSetAttribute. The codebase pattern is to wrap CUDA API calls with NVTE_CHECK_CUDA() for proper error handling.

Suggested change
cudaFuncSetAttribute(fused_topk_with_score_function_forward_kernel<DataType, BiasType>,
cudaFuncAttributeMaxDynamicSharedMemorySize, shared_memory_size);
NVTE_CHECK_CUDA(cudaFuncSetAttribute(fused_topk_with_score_function_forward_kernel<DataType, BiasType>,
cudaFuncAttributeMaxDynamicSharedMemorySize, shared_memory_size));

fused_topk_with_score_function_forward_kernel<DataType, BiasType>
<<<grid_size, kThreadsPerBlock, shared_memory_size, stream>>>(
logits, num_tokens, num_experts, topk, use_pre_softmax, num_groups, group_topk,
Expand Down Expand Up @@ -444,6 +446,8 @@ void fused_topk_with_score_function_backward_kernel_launcher(
num_experts * num_token_per_block * sizeof(DataType) // act_from_fwd
+ num_experts * num_token_per_block * sizeof(DataType) // comp_buf
+ num_experts * num_token_per_block * sizeof(bool); // routing_map
cudaFuncSetAttribute(fused_topk_with_score_function_backward_kernel<DataType>,
cudaFuncAttributeMaxDynamicSharedMemorySize, shared_memory_size);
Comment on lines +449 to +450
Copy link
Contributor

@greptile-apps greptile-apps bot Jan 9, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Following the pattern used elsewhere in the codebase (e.g., ln_fwd_cuda_kernel.cu), this call should be conditional:

Suggested change
cudaFuncSetAttribute(fused_topk_with_score_function_backward_kernel<DataType>,
cudaFuncAttributeMaxDynamicSharedMemorySize, shared_memory_size);
if (shared_memory_size >= 48 * 1024) {
cudaFuncSetAttribute(fused_topk_with_score_function_backward_kernel<DataType>,
cudaFuncAttributeMaxDynamicSharedMemorySize, shared_memory_size);
}

This attribute only needs to be set when exceeding the 48KB default limit.

Comment on lines +449 to +450
Copy link
Contributor

@greptile-apps greptile-apps bot Jan 9, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Missing error handling for cudaFuncSetAttribute. The codebase pattern is to wrap CUDA API calls with NVTE_CHECK_CUDA() for proper error handling.

Suggested change
cudaFuncSetAttribute(fused_topk_with_score_function_backward_kernel<DataType>,
cudaFuncAttributeMaxDynamicSharedMemorySize, shared_memory_size);
NVTE_CHECK_CUDA(cudaFuncSetAttribute(fused_topk_with_score_function_backward_kernel<DataType>,
cudaFuncAttributeMaxDynamicSharedMemorySize, shared_memory_size));

Note: If this suggestion doesn't match your team's coding style, reply to this and let me know. I'll remember it for next time!

fused_topk_with_score_function_backward_kernel<DataType>
<<<grid_size, kThreadsPerBlock, shared_memory_size, stream>>>(
routing_map, intermediate_output, grad_probs, num_tokens, num_experts, topk,
Expand Down