Skip to content

[https://nvbugs/5814309][fix] Use NCCL as fallback to avoid crash due to insufficient memory #10928

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

Merged
hyukn merged 2 commits into from
Jan 23, 2026
Merged

[https://nvbugs/5814309][fix] Use NCCL as fallback to avoid crash due to insufficient memory #10928

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
dfcec55
[https://nvbugs/5814309][fix] Use NCCL as fallback to avoid crash due…
hyukn Jan 22, 2026
91c49b7
remove unnecessary noexcept
hyukn Jan 22, 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 cpp/tensorrt_llm/common/customAllReduceUtils.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -115,7 +115,7 @@ inline AllReduceStrategyType selectStrategyLookUpTable(
|| num_token_index
>= AllReduceBestStrategyTable.at(sm_version).at(tp_index).at(fusion_op_index).at(hidden_size_index).size())
{
return AllReduceStrategyType::NCCL_SYMMETRIC;
return AllReduceStrategyType::NCCL;
}

return static_cast<AllReduceStrategyType>(
Expand Down
12 changes: 6 additions & 6 deletions cpp/tensorrt_llm/thop/allreduceOp.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -281,7 +281,7 @@ class AllreduceOp
std::vector<torch::Tensor> run(torch::Tensor const& input, torch::optional<torch::Tensor> const& residual,
torch::optional<torch::Tensor> const& norm_weight, torch::optional<torch::Tensor> const& scale,
torch::optional<torch::Tensor> const& bias, bool trigger_completion_at_end,
torch::optional<torch::Tensor> workspace) noexcept
torch::optional<torch::Tensor> workspace)
{
size_t size = input.numel();
size_t seq_len = input.size(0);
Expand Down Expand Up @@ -564,7 +564,7 @@ class AllreduceOp

std::vector<torch::Tensor> runLowPrecisionAllReduce(torch::Tensor const& input,
torch::optional<torch::Tensor> const& residual, torch::optional<torch::Tensor> const& norm_weight,
torch::optional<torch::Tensor> const& scale, torch::optional<torch::Tensor> const& bias) noexcept
torch::optional<torch::Tensor> const& scale, torch::optional<torch::Tensor> const& bias)
{
#ifdef ENABLE_FP8
auto stream = at::cuda::getCurrentCUDAStream(input.get_device());
Expand Down Expand Up @@ -632,8 +632,7 @@ class AllreduceOp
std::vector<torch::Tensor> runFusionAllReduce(torch::Tensor const& input,
torch::optional<torch::Tensor> const& residual, torch::optional<torch::Tensor> const& norm_weight,
torch::optional<torch::Tensor> const& scale, torch::optional<torch::Tensor> const& bias,
bool trigger_completion_at_end, torch::optional<torch::Tensor> workspace,
AllReduceStrategyType strategy) noexcept
bool trigger_completion_at_end, torch::optional<torch::Tensor> workspace, AllReduceStrategyType strategy)
{
// Should handle only Lamport implementation
auto stream = at::cuda::getCurrentCUDAStream(input.get_device());
Expand Down Expand Up @@ -1206,7 +1205,7 @@ class AllreduceOp

if (ifFallbackToNCCL(seq_len, message_size_bytes, max_workspace_size))
{
return AllReduceStrategyType::NCCL_SYMMETRIC;
return AllReduceStrategyType::NCCL;
Comment thread
hyukn marked this conversation as resolved.
}

// This rule based heuristic only chooses between NCCL_SYMMETRIC and MIN_LATENCY strategies.
Expand All @@ -1232,7 +1231,8 @@ class AllreduceOp

bool ifFallbackToNCCL(size_t seq_len, size_t message_size_bytes, size_t max_workspace_size)
{
// If messageSize is greater than maxWorkspaceSize or topology is unsuitable, use NCCL_SYMMETRIC fallback.
// If messageSize is greater than maxWorkspaceSize or topology is unsuitable, use NCCL fallback.
// TODO: Use NCCL_SYMMETRIC once the memory allocation issue is resolved.
if (message_size_bytes > max_workspace_size || !mIsP2PSupported || !mIsNVLINKSupported)
{
return true;
Expand Down
Loading