Repository navigation
[https://nvbugs/6899677][fix] Default NCCL_GRAPH_REGISTER=0 before creating the pipeline-parallel NCCL communicator - #20805
Conversation
|
No actionable comments were generated in the recent review.
 |
| Layer / File(s) | Summary |
|---|---|
NCCL environment setup cpp/tensorrt_llm/runtime/ncclCommunicator.cpp |
createComm sets NCCL_GRAPH_REGISTER to "0" when it is unset on Windows. On other platforms, it attempts to set the variable without overwriting an existing value. The existing NCCL_RUNTIME_CONNECT setup remains unchanged. |
Priority:
Estimated code review effort: 2 (Simple) | ~8 minutes
Change: Bug fix
Suggested labels: CUDA Graph, api-compatible
Merge Risk: 
31c97
-
The change preserves explicit environment settings. No merge-blocking issue was established.
Pre-merge checks: 
-
Check name Status Explanation Resolution Docstring Coverage Docstring coverage is 0.00% which is insufficient. The required threshold is 80.00%. Docstring coverage is scoped to functions touched by this diff with cyclomatic complexity of at least 2. Documentat… Write docstrings for the functions missing them to satisfy the coverage threshold. -
Check name Status Explanation Title check The title follows the required NVBugs and fix format and clearly identifies the main change: setting NCCL_GRAPH_REGISTER=0 before creating the pipeline-parallel communicator. Description check The description clearly explains the failure, root cause, implementation, behavior change, limitations, and extensive test evidence. It includes the required Description, Test Coverage, and PR Checkli… Linked Issues check Check skipped because no linked issues were found for this pull request. Out of Scope Changes check Check skipped because no linked issues were found for this pull request. Full details: Description check
-
Explanation
The description clearly explains the failure, root cause, implementation, behavior change, limitations, and extensive test evidence. It includes the required Description, Test Coverage, and PR Checklist sections.
-
Full details: Docstring Coverage
-
Explanation
Docstring coverage is 0.00% which is insufficient. The required threshold is 80.00%. Docstring coverage is scoped to functions touched by this diff with cyclomatic complexity of at least 2. Documentation-only edits and deletions within surviving functions are included; unchanged diff context and fully deleted functions are excluded. Analyzed 1 functions across 1 files.
Resolution
Write docstrings for the functions missing them to satisfy the coverage threshold.
-
Finishing Touches
-
🧪 Generate unit tests (beta)
-
- Create a new PR
-
- Autofix · Keep fixing CodeRabbit findings and required CI, and resolving merge conflicts
Comment @coderabbitai help to get the list of available commands.
|
Bug engineer ownership @YihuiLu512, as the engineer assigned to the linked bug, you are responsible for driving this bot PR to resolution. Please coordinate review, address feedback and required code changes, get CI passing, and verify the original failure before driving it to merge. If the PR is obsolete or superseded, close it with a reason. Keep this PR updated with blockers and the next step. |
|
Automatically added "ci: full pre-merge approved" because this PR has satisfied the required GitHub review approvals. Unresolved review conversations and other required checks remain independent merge requirements. |
…eating the pipeline-parallel NCCL communicator getComm() sets NCCL_GRAPH_REGISTER=0, without overwriting a user value, before it creates a communicator. NcclCommunicator::createComm, which creates the pipeline-parallel communicator used by PPCommNCCL, did not. With default settings, NCCL caches the value on its first read, which happens for an operation enqueued during CUDA graph capture. Where that first read came before any getComm() communicator existed, the process kept the default (1). This happened in MPI worker processes that had served a pipeline-parallel LLM without tensor or context parallelism whose send/recv was captured into CUDA graphs: a tensor- and pipeline-parallel LLM that they served later captured with graph registration enabled and hung. Apply the same default in createComm. Signed-off-by: Yihui Lu <269394165+YihuiLu512@users.noreply.github.com>
d6e2b09 to
3b395a8
Compare
3b395a8 to
31c9765
Compare
YihuiLu512
changed the title
|
/bot run --disable-fail-fast |
|
PR_Github #81395 [ run ] triggered by Bot. Commit: |
|
PR_Github #81395 [ run ] completed with state
|
|
/bot run --disable-fail-fast |
|
PR_Github #81570 [ run ] triggered by Bot. Commit: |
|
PR_Github #81570 [ run ] completed with state |
Description
TestNemotronV3Ultra::test_nvfp4_parallelism[TEP2_PP2](T; tp2/ep2/pp2) hangs during executor creation when it runs inMPI worker processes that previously served
TestDeepSeekV3Lite::test_nvfp4_4gpus[moe_backend=CUTLASS-mtp_nextn=2-pp4-fp8kv=True-attention_dp=True-cuda_graph=True-overlap_scheduler=True-low_precision_combine=False-torch_compile=False](the pp4 test); the accuracy tests reuse the 4-worker MPI pool between tests. All ranks stop in
torch.cuda.synchronize()right after capturing the firstCUDA graph (batch size 32), and the GPUs spin at 100% utilization.
Root cause:
getComm()(cpp/tensorrt_llm/common/opUtils.cpp) setsNCCL_GRAPH_REGISTER=0, without overwriting auser value, before it creates a communicator ("Disable graph register to avoid startup hangs").
NcclCommunicator::createComm, which creates the pipeline-parallel communicator used byPPCommNCCL, does not. Withdefault NCCL settings, NCCL reads
NCCL_GRAPH_REGISTERonly for operations enqueued during CUDA graph capture, and itcaches the first value it reads for the rest of the process. The pp4 test has no tensor or context parallelism, so it
creates no
getComm()communicator, but it captures its pipeline send/recv into CUDA graphs; NCCL therefore caches thedefault, 1. T then runs in the same worker processes and captures with graph registration enabled, even though its own
getComm()sets the variable to 0, and hangs.Fix: apply the same
NCCL_GRAPH_REGISTER=0default inNcclCommunicator::createComm, next to theNCCL_RUNTIME_CONNECTdefault it already sets, so the value is in place before the pipeline-parallel communicator iscreated.
Behavior change: only processes that use
PPCommNCCL(MPI mode) set the new default; it is process-wide, and childprocesses started afterwards inherit it. With default NCCL settings, behavior changes only in processes where
createCommnow sets the variable before NCCL first reads it and nogetComm()communicator existed by then: theirCUDA graph captures of NCCL operations used to run with graph registration on and now run with it off. The performance
effect of this was not measured. A value set in the environment is still honored; a value set only in
nccl.confcan bepre-empted, as it already can be by
getComm().This PR was first opened against release/1.3, where QA hit the hang. NVBug 6899677 now targets 1.4.0, so the PR now
targets main.
Evidence (B300; the CI image of release/1.3 351f9eb, the commit of the failing CI run, with NCCL 2.30.7, and a wheel
rebuilt from that commit; 4 GPUs unless noted; counts are runs):
or experimental patches, or with NCCL debug logging, also hang, 7/7.
NCCL_GRAPH_REGISTER=0exported for the whole job: passes, 3/3.PPCommNCCLsets the default before creating its communicator):passes, 3/3.
NCCL_GRAPH_REGISTER=1exported: hangs at the same point, 3/3. T alone without it:passes, 3/3 (one of these runs used 8 GPUs; the other two carry instrumentation or NCCL debug logging).
NCCL_DEBUG_SUBSYS=INIT,ENV,REG): in a hanging pp4 -> T run, no process logsNCCL_GRAPH_REGISTER set by environment, and each process logsfailed to IPC register userbufflines, both duringthe pp4 test (264-528 per process) and during T (289-291 per process). In a passing T-alone run, every process logs
NCCL_GRAPH_REGISTER set by environment to 0.after itsgetComm()communicator is created, and none of thosefailures.
tp4in place ofpp4(otherwise the same parameters) before T: T passes(1 run).
ncclParamGraphRegister()is guarded by CUDA graph capture, exceptone in
transport/generic.ccthat the defaultNCCL_LOCAL_REGISTER=1short-circuits;NCCL_PARAMvalues are cachedafter the first read unless listed in
NCCL_NO_CACHE.Test Coverage
T is in the QA list only (
tests/integration/test_lists/qa/llm_function_core.txt), and the hang needs an earlier testin the same MPI worker processes, such as the pp4 test, so the tests were run directly on B300.
The runs below are on release/1.3, not on main. The
9671b46b45builds apply this PR's change unmodified; the351f9ebe02builds apply the same code change with a differently worded code comment. The wheels with the fix werebuilt for sm103 only. main and both release/1.3
commits below have identical
cpp/tensorrt_llm/runtime/ncclCommunicator.cpp(before this change),cpp/tensorrt_llm/common/opUtils.cpp, andPPCommNCCL/init_pp_commintensorrt_llm/_torch/distributed/communicator.py; main pins the same NCCL version as9671b46b45(2.31.2, indocker/common/install_cuda_libs.sh).On the B300 platform of the failing CI run (driver 580.126.09; 8 GPUs per job), with four tests from the failing CI run
in the order they ran there:
TestDeepSeekV3Lite::test_nvfp4_4gpus[moe_backend=CUTLASS-mtp_nextn=0-ep4-fp8kv=True-attention_dp=True-cuda_graph=True-overlap_scheduler=True-low_precision_combine=False-torch_compile=False](the ep4 test), the pp4 test,
TestDeepSeekV4ProDSpark::test_gsm8k_dep8_megamoe_deepgemm(on main:TestDeepSeekV4ProDSpark::test_gsm8k_megamoe_deepgemm[dp8]), then T, which reuses the 4 MPI workers of the ep4 andpp4 tests:
9671b46b45+ this fix, on the CI image of9671b46b45(NCCL 2.31.2), with a wheel built from thatsource:
NCCL_GRAPH_REGISTER=1exported: T hangs at the same point, 1 run.351f9ebe02+ this fix, on the CI image of351f9ebe02(NCCL 2.30.7), with a wheel built from thatsource:
351f9ebe02using the build command of thefailing CI run: T hangs, 13/13, and in every run its log stops at the first CUDA graph capture (batch size 32).
NCCL_GRAPH_REGISTER=1exported: T hangs at the same point, 1 run.On other B300 machines (4 GPUs per job), without the ep4 and DeepSeek-V4 tests:
9671b46b45+ this fix, on the CI image of9671b46b45(NCCL 2.31.2), with a wheel built from thatsource:
NCCL_GRAPH_REGISTER=1exported: T hangs at the same point, 1 run. The fix does not override avalue set in the environment, so this shows that with the fix the hang still depends on this variable.
351f9ebe02+ this fix, in the setup of the Evidence section, with a wheel built from that source:NCCL_GRAPH_REGISTER=1exported: T hangs at the same point, 1 run.The repair bot that first opened this PR also reported, on B300 at release/1.3
0f37627b2f, thatTestNemotronV3Ultra::test_nvfp4_parallelism[TP2_PP2]hangs the same way after a DeepSeek-V3-Lite nvfp4 pp4 test, andpasses with this fix.
PR Checklist
Dev Engineer Review
NcclCommunicator::createCommnow setsNCCL_GRAPH_REGISTER=0before creating the communicator. It preserves an existing value on Windows and other platforms. This aligns the default withgetComm()and addresses cases where NCCL reads the setting during an earlier CUDA graph capture. No public API changes are reported. Test results and review findings were not supplied.QA Engineer Review
No test changes.
Per-File QA Perspective
cpp/tensorrt_llm/runtime/ncclCommunicator.cpp: Verify that communicator creation sets the default before NCCL reads the setting, and that an explicitly set value remains unchanged on all supported platforms.