[#12715][fix] disable NCCL_SYMMETRIC tactic on GB10 (DGX Spark) by nv-lschneider · Pull Request #12902 · NVIDIA/TensorRT-LLM

nv-lschneider · 2026-04-09T20:37:00Z

Summary by CodeRabbit

Bug Fixes
- Fixed compatibility issues with GB10 CUDA devices in distributed reduction operations.
- Improved fallback behavior when optimization cache is unavailable.
Performance
- Optimized distributed reduction tactics for specific hardware platforms.

Description

This PR disables NCCL symmetric on spark compute notes.
#12715 reports that there is a failure on that platform.
Until the issue can be resolved from within NCCL this work around should enable the platform again without impacting other platforms.
We are investigating actively this failure in NCCL and will resolve it, when that happens we can remove this WAR again.

Test Coverage

I do not have access to a dual spark machine setup, so I need help with testing and verifying this fix.

PR Checklist

Please review the following before submitting your PR:

PR description clearly explains what and why. If using CodeRabbit's summary, please make sure it makes sense.
PR Follows TRT-LLM CODING GUIDELINES to the best of your knowledge.
Test cases are provided for new code paths (see test instructions)
Any new dependencies have been scanned for license and vulnerabilities
CODEOWNERS updated if ownership changes
Documentation updated as needed
Update tava architecture diagram if there is a significant design change in PR.
The reviewers assigned automatically/manually are appropriate for the PR.
Please check this after reviewing the above items as appropriate for this PR.

GitHub Bot Help

To see a list of available CI bot commands, please comment /bot help.

coderabbitai · 2026-04-09T20:43:38Z

📝 Walkthrough

Walkthrough

Added a cached GPU platform detection helper to identify GB10 devices. Modified AllReduceRunner to exclude the NCCL_SYMMETRIC strategy on GB10 GPUs while keeping both NCCL and NCCL_SYMMETRIC available on other devices. Adjusted fallback strategy selection accordingly.

Changes

Cohort / File(s)	Summary
AllReduce GPU Compatibility `tensorrt_llm/_torch/custom_ops/torch_custom_ops.py`	Added `_is_gb10()` platform detection helper; modified `AllReduceRunner.get_valid_tactics()` to exclude `NCCL_SYMMETRIC` on GB10; updated `AllReduceRunner.forward()` fallback to select `NCCL` on GB10 and `NCCL_SYMMETRIC` elsewhere.

Estimated code review effort

🎯 2 (Simple) | ⏱️ ~10 minutes

🚥 Pre-merge checks | ✅ 1 | ❌ 2

❌ Failed checks (1 warning, 1 inconclusive)

Check name	Status	Explanation	Resolution
Docstring Coverage	⚠️ Warning	Docstring coverage is 25.00% which is insufficient. The required threshold is 80.00%.	Write docstrings for the functions missing them to satisfy the coverage threshold.
Description check	❓ Inconclusive	The description explains the issue and workaround, but lacks formal test coverage details and does not address all checklist items adequately.	Clarify what tests will verify this GB10-specific behavior and confirm the author's understanding of testing requirements for platform-specific changes.

✅ Passed checks (1 passed)

Check name	Status	Explanation
Title check	✅ Passed	The title accurately summarizes the main change: disabling NCCL_SYMMETRIC tactic on GB10 (DGX Spark), references the issue (`#12715`), and follows the required format with type [fix].

_{✏️ Tip: You can configure your own custom pre-merge checks in the settings.}

✨ Finishing Touches

🧪 Generate unit tests (beta)

Create PR with unit tests

_{Comment @coderabbitai help to get the list of available commands and usage tips.}

coderabbitai

Actionable comments posted: 1

🤖 Prompt for all review comments with AI agents

Verify each finding against the current code and only fix it if needed.

Inline comments:
In `@tensorrt_llm/_torch/custom_ops/torch_custom_ops.py`:
- Around line 33-37: _is_gb10 currently calls torch.cuda.get_device_name() with
no device and caches a single global result, which can be wrong if the current
CUDA device changes or if CUDA is unavailable. Change _is_gb10 to accept an
optional device argument (e.g., def _is_gb10(device: Optional[int] = None) ->
bool), first check torch.cuda.is_available() and return False when CUDA is
unavailable, resolve device to torch.cuda.current_device() if None, call
torch.cuda.get_device_name(device) with that device id, and keep `@lru_cache` so
results are cached per-device; update call sites to pass the device when they
have a specific CUDA device context (or call with no arg to auto-resolve the
current device).

🪄 Autofix (Beta)

Fix all unresolved CodeRabbit comments on this PR:

Push a commit to this branch (recommended)
Create a new PR with the fixes

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Pro

Run ID: 3f68adeb-1c67-4833-a08c-0d6d163cff41

📥 Commits

Reviewing files that changed from the base of the PR and between aed6b6a and df9413f.

📒 Files selected for processing (1)

tensorrt_llm/_torch/custom_ops/torch_custom_ops.py

Tabrizian · 2026-04-09T20:48:10Z

Thanks for the quick fix!

nv-lschneider · 2026-04-09T20:49:10Z

/bot run --disable-fail-fast

tensorrt-cicd · 2026-04-09T20:55:04Z

PR_Github #42581 [ run ] triggered by Bot. Commit: df9413f Link to invocation

tensorrt-cicd · 2026-04-10T16:07:32Z

PR_Github #42581 [ run ] completed with state SUCCESS. Commit: df9413f
/LLM/main/L0_MergeRequest_PR pipeline #33309 completed with status: 'FAILURE'

CI Report

⚠️ Action Required:

Please check the failed tests and fix your PR
If you cannot view the failures, ask the CI triggerer to share details
Once fixed, request an NVIDIA team member to trigger CI again

Link to invocation

Tabrizian · 2026-04-10T16:13:41Z

/bot run --disable-fail-fast

nv-lschneider · 2026-04-10T16:13:48Z

/bot run --disable-fail-fast

tensorrt-cicd · 2026-04-10T16:20:38Z

PR_Github #42715 [ run ] triggered by Bot. Commit: df9413f Link to invocation

tensorrt-cicd · 2026-04-10T16:20:52Z

PR_Github #42716 [ run ] triggered by Bot. Commit: df9413f Link to invocation

tensorrt-cicd · 2026-04-10T16:20:55Z

PR_Github #42715 [ run ] completed with state ABORTED. Commit: df9413f

Link to invocation

tensorrt-cicd · 2026-04-10T23:06:27Z

PR_Github #42716 [ run ] completed with state SUCCESS. Commit: df9413f
/LLM/main/L0_MergeRequest_PR pipeline #33406 completed with status: 'SUCCESS'

CI Report

Link to invocation

theherrovn · 2026-06-03T12:41:08Z

@nv-lschneider — I've been investigating the remaining hang after this PR's WAR fix.

Root Cause Analysis

The WAR correctly disables NCCL_SYMMETRIC on GB10, which fixes the segfault. But there's a second bug in the same initialization path:

MNNVL Fabric Check Still Runs

ncclMnnvlCheck() (NCCL's src/mnnvl.cc) is called during communicator init when nNodes > 1. On dual DGX Spark:

NVML reports fabricInfo.state in an unexpected state (not COMPLETED, not NOT_SUPPORTED)
ncclMnnvlCheck() proceeds past Step 3 into Step 5 — FABRIC handle export/import
cuMemExportToShareableHandle(CU_MEM_HANDLE_TYPE_FABRIC) fails/hangs on GB10 hardware with no actual NVLink fabric

Evidence

The remaining hang occurs at the "checking fabric" step — after NCCL_SYMMETRIC is disabled, but before the autotuner warmup completes. This is exactly where ncclMnnvlCheck() runs.

Proposed Fix

Extend this PR to also set NCCL_MNNVL_ENABLE=0 on GB10:

@lru_cache(maxsize=None)
def _is_gb10() -> bool:
    """Return True on GB10 (DGX Spark). NCCL_SYMMETRIC and MNNVL fabric are unsupported."""
    return "GB10" in torch.cuda.get_device_name()

# In initialization:
if _is_gb10():
    os.environ["NCCL_MNNVL_ENABLE"] = "0"  # No NVLink fabric on GB10

NCCL_MNNVL_ENABLE=0 cleanly skips the entire ncclMnnvlCheck() path (NCCL init.cc:1024), preventing the fabric handle exchange that hangs on GB10.

Verification

Workaround confirmed by setting NCCL_MNNVL_ENABLE=0 manually on any dual DGX Spark setup
NCCL_MNNVL_ENABLE=2 is the default (auto); setting to 0 forces skip
Works independently of NCCL_SYMMETRIC — they control different code paths

Happy to add this to the PR if it makes sense.

nv-lschneider · 2026-06-03T20:01:13Z

@nv-lschneider — I've been investigating the remaining hang after this PR's WAR fix.

Root Cause Analysis

The WAR correctly disables NCCL_SYMMETRIC on GB10, which fixes the segfault. But there's a second bug in the same initialization path:

MNNVL Fabric Check Still Runs

ncclMnnvlCheck() (NCCL's src/mnnvl.cc) is called during communicator init when nNodes > 1. On dual DGX Spark:

NVML reports fabricInfo.state in an unexpected state (not COMPLETED, not NOT_SUPPORTED)

ncclMnnvlCheck() proceeds past Step 3 into Step 5 — FABRIC handle export/import

cuMemExportToShareableHandle(CU_MEM_HANDLE_TYPE_FABRIC) fails/hangs on GB10 hardware with no actual NVLink fabric

Evidence

The remaining hang occurs at the "checking fabric" step — after NCCL_SYMMETRIC is disabled, but before the autotuner warmup completes. This is exactly where ncclMnnvlCheck() runs.

Proposed Fix

Extend this PR to also set NCCL_MNNVL_ENABLE=0 on GB10:
@lru_cache(maxsize=None)
def _is_gb10() -> bool:
    """Return True on GB10 (DGX Spark). NCCL_SYMMETRIC and MNNVL fabric are unsupported."""
    return "GB10" in torch.cuda.get_device_name()

# In initialization:
if _is_gb10():
    os.environ["NCCL_MNNVL_ENABLE"] = "0"  # No NVLink fabric on GB10
NCCL_MNNVL_ENABLE=0 cleanly skips the entire ncclMnnvlCheck() path (NCCL init.cc:1024), preventing the fabric handle exchange that hangs on GB10.

Verification

Workaround confirmed by setting NCCL_MNNVL_ENABLE=0 manually on any dual DGX Spark setup

NCCL_MNNVL_ENABLE=2 is the default (auto); setting to 0 forces skip

Works independently of NCCL_SYMMETRIC — they control different code paths

Happy to add this to the PR if it makes sense.

Thanks for taking a deeper look into this. This helps quite a lot.
I am unsure about if it is independent of NCCL_SYMMETRIC, since NCCL_SYMMETRIC has an option to run with MNNVL, so I suspect that they are related.
And possibly with NCCL_MNNVL_ENABLE=0 we can run NCCL_SYMMETRIC just as before.
If that is the case, we should replace this with the current fix.

I will also take the AI and check if this is a NCCL version specific problem, such that it can be resolved in future NCCL versions.
Thank you for your analysis.

Feel free to add your commits to my existing PR.

nv-lschneider · 2026-06-04T22:34:33Z

I tested the findings from @theherrovn to use the MNNVL disable switch and it works together with disabling NCCL_SYMMETRIC.

Additionally, I verified that this is fixed in NCCL 2.30.4 and should be automatically fixed with the next update of the NCCL version.

This PR submits a work around to set the env var and disables NCCL_SYMMETRIC on gb10 environments.

This has been tested on a dual-spark setup.
With the RC10 release container it hangs, patching the container with these changes, resolves the hang and restores normal allreduce behavior.

nv-lschneider · 2026-06-04T22:39:48Z

/bot run

tensorrt-cicd · 2026-06-04T22:45:12Z

PR_Github #52174 [ run ] triggered by Bot. Commit: ef06316 Link to invocation

tensorrt-cicd · 2026-06-04T23:52:14Z

PR_Github #52174 [ run ] completed with state SUCCESS. Commit: ef06316
/LLM/main/L0_MergeRequest_PR pipeline #41495 completed with status: 'FAILURE'

CI Report

⚠️ Action Required:

Please check the failed tests and fix your PR
If you cannot view the failures, ask the CI triggerer to share details
Once fixed, request an NVIDIA team member to trigger CI again

CI Agent Failure Analysis

Link to invocation

theherrovn · 2026-06-05T05:27:11Z

@nv-lschneider Great to hear the MNNVL disable switch works alongside the NCCL_SYMMETRIC fix on real dual-Spark hardware. And good to know NCCL 2.30.4 resolves this permanently — the WAR is the right approach until the upgrade lands.

Thanks for testing and confirming!

tensorrt-cicd · 2026-06-10T22:12:03Z

PR_Github #53382 [ run ] triggered by Bot. Commit: 80f91a2 Link to invocation

tensorrt-cicd · 2026-06-11T02:53:11Z

PR_Github #53382 [ run ] completed with state FAILURE. Commit: 80f91a2
/LLM/main/L0_MergeRequest_PR pipeline #42559 completed with status: 'FAILURE'

CI Report

⚠️ Action Required:

Please check the failed tests and fix your PR
If you cannot view the failures, ask the CI triggerer to share details
Once fixed, request an NVIDIA team member to trigger CI again

CI Agent Failure Analysis

Link to invocation

nv-lschneider · 2026-06-11T12:41:53Z

/bot run

tensorrt-cicd · 2026-06-11T12:47:22Z

PR_Github #53569 [ run ] triggered by Bot. Commit: d237242 Link to invocation

tensorrt-cicd · 2026-06-11T13:29:07Z

PR_Github #53569 [ run ] completed with state FAILURE. Commit: d237242
/LLM/main/L0_MergeRequest_PR pipeline #42718 completed with status: 'FAILURE'

CI Report

⚠️ Action Required:

Please check the failed tests and fix your PR
If you cannot view the failures, ask the CI triggerer to share details
Once fixed, request an NVIDIA team member to trigger CI again

CI Agent Failure Analysis

Link to invocation

nv-lschneider · 2026-06-11T16:14:53Z

/bot run

tensorrt-cicd · 2026-06-11T16:21:01Z

PR_Github #53611 [ run ] triggered by Bot. Commit: 5b55103 Link to invocation

tensorrt-cicd · 2026-06-11T17:00:15Z

PR_Github #53611 [ run ] completed with state FAILURE. Commit: 5b55103
/LLM/main/L0_MergeRequest_PR pipeline #42757 completed with status: 'FAILURE'

CI Report

⚠️ Action Required:

Please check the failed tests and fix your PR
If you cannot view the failures, ask the CI triggerer to share details
Once fixed, request an NVIDIA team member to trigger CI again

CI Agent Failure Analysis

Link to invocation

Signed-off-by: Ludwig Schneider <lschneider@nvidia.com>

NCCL_SYMMETRIC and NCCL MNNVL need to be disabled on GB10 when running affected NCCL runtimes. Query ncclGetVersion at runtime and only apply the autotuner filter plus NCCL_MNNVL_ENABLE=0 export for NCCL versions older than 2.30.4. NCCL already fixed this in 2.30.4 and newer, so upgraded installations keep NCCL_SYMMETRIC in the tactic set and do not receive the MNNVL env workaround. Credit to theherrovn for investigating the GB10 hang and identifying the NCCL_MNNVL_ENABLE=0 env export. Investigated-by: theherrovn Signed-off-by: Ludwig Schneider <lschneider@nvidia.com>

nv-lschneider · 2026-06-11T20:08:11Z

/bot run

tensorrt-cicd · 2026-06-11T20:14:45Z

PR_Github #53659 [ run ] triggered by Bot. Commit: fdd8d33 Link to invocation

tensorrt-cicd · 2026-06-11T21:44:04Z

PR_Github #53659 [ run ] completed with state SUCCESS. Commit: fdd8d33
/LLM/main/L0_MergeRequest_PR pipeline #42801 completed with status: 'FAILURE'

CI Report

⚠️ Action Required:

Please check the failed tests and fix your PR
If you cannot view the failures, ask the CI triggerer to share details
Once fixed, request an NVIDIA team member to trigger CI again

CI Agent Failure Analysis

Link to invocation

nv-lschneider · 2026-06-11T22:33:07Z

/bot run --disable-fail-fast

tensorrt-cicd · 2026-06-11T22:38:11Z

PR_Github #53699 [ run ] triggered by Bot. Commit: fdd8d33 Link to invocation

tensorrt-cicd · 2026-06-11T22:39:23Z

PR_Github #53700 [ run ] triggered by Bot. Commit: fdd8d33 Link to invocation

tensorrt-cicd · 2026-06-11T22:39:27Z

PR_Github #53699 [ run ] completed with state ABORTED. Commit: fdd8d33

Link to invocation

tensorrt-cicd · 2026-06-12T04:57:36Z

PR_Github #53700 [ run ] completed with state FAILURE. Commit: fdd8d33
/LLM/main/L0_MergeRequest_PR pipeline #42832 completed with status: 'FAILURE'

CI Report

⚠️ Action Required:

Please check the failed tests and fix your PR
If you cannot view the failures, ask the CI triggerer to share details
Once fixed, request an NVIDIA team member to trigger CI again

CI Agent Failure Analysis

Link to invocation

nv-lschneider · 2026-06-12T16:38:45Z

/bot run --disable-fail-fast

tensorrt-cicd · 2026-06-12T16:44:53Z

PR_Github #53917 [ run ] triggered by Bot. Commit: fdd8d33 Link to invocation

tensorrt-cicd · 2026-06-12T20:25:25Z

PR_Github #53917 [ run ] completed with state SUCCESS. Commit: fdd8d33
/LLM/main/L0_MergeRequest_PR pipeline #43012 completed with status: 'SUCCESS'
Pipeline passed with automatic retried tests. Check the rerun report for details.

CI Report

Link to invocation

nv-lschneider requested a review from a team as a code owner April 9, 2026 20:37

nv-lschneider requested a review from hyukn April 9, 2026 20:37

github-actions Bot assigned nv-lschneider Apr 9, 2026

nv-lschneider requested a review from Tabrizian April 9, 2026 20:38

coderabbitai Bot reviewed Apr 9, 2026

View reviewed changes

Comment thread tensorrt_llm/_torch/custom_ops/torch_custom_ops.py

Tabrizian approved these changes Apr 9, 2026

View reviewed changes

nv-lschneider mentioned this pull request Apr 10, 2026

[Bug]: Multi-node regression: segfault/hang in NCCL fabric detection on dual DGX Spark (1.3.0rc10 vs 1.2.0rc6) #12715

Closed

5 tasks

nv-lschneider force-pushed the lschneider/war-dual-dgx-nccl-symmetric branch 2 times, most recently from f66399e to 80f126d Compare June 4, 2026 22:14

nv-lschneider requested a review from a team as a code owner June 4, 2026 22:14

nv-lschneider requested a review from HuiGao-NV June 4, 2026 22:14

nv-lschneider force-pushed the lschneider/war-dual-dgx-nccl-symmetric branch from 80f126d to ef06316 Compare June 4, 2026 22:37

nv-lschneider force-pushed the lschneider/war-dual-dgx-nccl-symmetric branch from 80f91a2 to d237242 Compare June 11, 2026 12:41

nv-lschneider force-pushed the lschneider/war-dual-dgx-nccl-symmetric branch from d237242 to 5b55103 Compare June 11, 2026 16:14

nv-lschneider added 2 commits June 11, 2026 15:07

[None][bugfix] disable NCCL_SYMMETRIC tactic on GB10 (DGX Spark)

907c86a

Signed-off-by: Ludwig Schneider <lschneider@nvidia.com>

nv-lschneider force-pushed the lschneider/war-dual-dgx-nccl-symmetric branch from 5b55103 to fdd8d33 Compare June 11, 2026 20:07

nv-lschneider requested a review from Tabrizian June 12, 2026 20:27

Tabrizian approved these changes Jun 12, 2026

View reviewed changes

hyukn approved these changes Jun 13, 2026

View reviewed changes

hyukn merged commit cd65070 into NVIDIA:main Jun 13, 2026
7 checks passed

coderabbitai Bot mentioned this pull request Jun 18, 2026

[https://nvbugs/6322073][fix] Add _needs_x86_nccl_pp_workaround() and init_nccl_pp_workaround() helpers… #15480

Closed

2 tasks

nv-lschneider mentioned this pull request Jun 23, 2026

[#12715][fix] Disable NCCL window buffers on GB10 #15559

Merged

1 task

Uh oh!

Conversation

nv-lschneider commented Apr 9, 2026 • edited by coderabbitai Bot Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary by CodeRabbit

Description

Test Coverage

PR Checklist

GitHub Bot Help

Uh oh!

coderabbitai Bot commented Apr 9, 2026

Walkthrough

Changes

Estimated code review effort

❌ Failed checks (1 warning, 1 inconclusive)

Uh oh!

coderabbitai Bot left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Tabrizian commented Apr 9, 2026

Uh oh!

nv-lschneider commented Apr 9, 2026

Uh oh!

tensorrt-cicd commented Apr 9, 2026

Uh oh!

tensorrt-cicd commented Apr 10, 2026

Uh oh!

Tabrizian commented Apr 10, 2026

Uh oh!

nv-lschneider commented Apr 10, 2026

Uh oh!

tensorrt-cicd commented Apr 10, 2026

Uh oh!

tensorrt-cicd commented Apr 10, 2026

Uh oh!

tensorrt-cicd commented Apr 10, 2026

Uh oh!

tensorrt-cicd commented Apr 10, 2026

Uh oh!

theherrovn commented Jun 3, 2026

Root Cause Analysis

MNNVL Fabric Check Still Runs

Evidence

Proposed Fix

Verification

Uh oh!

nv-lschneider commented Jun 3, 2026

Root Cause Analysis

MNNVL Fabric Check Still Runs

Evidence

Proposed Fix

Verification

Uh oh!

nv-lschneider commented Jun 4, 2026

Uh oh!

nv-lschneider commented Jun 4, 2026

Uh oh!

tensorrt-cicd commented Jun 4, 2026

Uh oh!

tensorrt-cicd commented Jun 4, 2026

Uh oh!

theherrovn commented Jun 5, 2026

Uh oh!

tensorrt-cicd commented Jun 10, 2026

Uh oh!

tensorrt-cicd commented Jun 11, 2026

Uh oh!

nv-lschneider commented Jun 11, 2026

Uh oh!

tensorrt-cicd commented Jun 11, 2026

Uh oh!

tensorrt-cicd commented Jun 11, 2026

Uh oh!

nv-lschneider commented Jun 11, 2026

Uh oh!

tensorrt-cicd commented Jun 11, 2026

Uh oh!

tensorrt-cicd commented Jun 11, 2026

Uh oh!

nv-lschneider commented Apr 9, 2026 •

edited by coderabbitai Bot

Loading