[https://nvbugs/6801108][fix] Set NCCL env vars before ncclGetUniqueId in getComm - #19596
Conversation
…d in getComm getComm() called setenv() for NCCL_RUNTIME_CONNECT, NCCL_GRAPH_REGISTER and NCCL_NVLS_ENABLE after getUniqueId(). On the group root, ncclGetUniqueId() starts NCCL's bootstrap root thread, and since NCCL 2.30.7 that thread calls getenv() the first time it uses ncclSocketDefaultMagic(). setenv() is not thread-safe: adding a variable can reallocate environ while the bootstrap thread is scanning it. The root rank then occasionally crashed with SIGSEGV and the remaining ranks waited forever in NCCL bootstrap, which showed up as a hang in the TTP cases of test_configurable_moe_multi_gpu. Move the environment updates before getUniqueId(), matching the order already used by NcclCommunicator::createComm(). Signed-off-by: xxi <xxi@nvidia.com>
|
Navigate logical layers of code changes, visualize relationships, and explore their blast radius. No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info⚙️ Run configurationConfiguration used: Repository: NVIDIA/TensorRT-LLM/.coderabbit.yaml Review profile: CHILL Plan: Enterprise Run ID: 📒 Files selected for processing (1)
Included review availability: Your plan provides up to 12 included reviews per hour; 11 remain after this review. Walkthrough
ChangesNCCL environment setup
Priority: ⚪ Not assessed Estimated code review effort: 2 (Simple) | ~5 minutes Change: Bug fix Suggested reviewers: Merge Risk: ⚪ Minimal · up to No actionable merge-blocking risk is evident from the supplied change; it is ready for normal merge checks. 🚥 Pre-merge checks | ✅ 5✅ Passed checks (5 passed)
✨ Finishing Touches🧪 Generate unit tests (beta)
Comment |
|
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. |
|
/bot run --disable-fail-fast |
|
PR_Github #75745 [ run ] triggered by Bot. Commit: |
|
PR_Github #75745 [ run ] completed with state
|
|
/bot run --disable-fail-fast |
|
PR_Github #75781 [ run ] triggered by Bot. Commit: |
|
PR_Github #75781 [ run ] completed with state
|
|
/bot run --disable-fail-fast |
|
PR_Github #75793 [ run ] triggered by Bot. Commit: |
|
PR_Github #75793 [ run ] completed with state
|
|
/bot run --disable-fail-fast |
|
PR_Github #75795 [ run ] triggered by Bot. Commit: |
|
Evaluate behavioral incompatibilities involving the supplied PR changes and the Verify the three full commit IDs and their merge-base. For divergent branches, When head contains target ( First identify changed or removed interfaces and symbols, then search repository Classify findings in prose as cross-branch interactions or PR-local compatibility Use read-only source and Git inspection. Do not modify checked-out files, execute Assess independent findings separately. Uncertainty about one path does not
Write exactly one standalone heading
Then give the verdict, findings or inspected paths, and limitations. For PASS and Copy the supplied identity and commit IDs verbatim. Always include the result Analyze only these fixed revisions in NVIDIA/TensorRT-LLM, not the hosting PR: |
|
Tip For best results, initiate chat on the files or code changes. 🧩 Analysis chain🏁 Script executed: #!/bin/bash
set -uo pipefail
head=20130077a0ba487470425a4af1a22b7b01ac9760
target=21d2f1c253c60da241c187f9489e882f29a31425
base=adea5c0720f266ef3ec8490b638498f491af682d
for rev in "$head" "$target" "$base"; do
git cat-file -t "$rev"
git rev-parse --verify "$rev^{commit}"
done
git merge-base --all "$head" "$target"
git diff --stat "$base" "$head"
git diff --stat "$base" "$target"
git diff --name-status "$base" "$head"
git diff --name-status "$base" "$target"
git diff "$base" "$head" -- cpp/tensorrt_llm/common/opUtils.cppLength of output: 111045 🏁 Script executed: #!/bin/bash
set -uo pipefail
head=20130077a0ba487470425a4af1a22b7b01ac9760
target=21d2f1c253c60da241c187f9489e882f29a31425
base=adea5c0720f266ef3ec8490b638498f491af682d
ast-grep outline cpp/tensorrt_llm/common/opUtils.cpp --items all
for rev in "$head" "$target" "$base"; do
echo "=== $rev opUtils ==="
git show "$rev:cpp/tensorrt_llm/common/opUtils.cpp" | nl -ba | sed -n '1,215p'
echo "=== $rev references ==="
git grep -n -E '\b(getComm|getUniqueId|ipcNvlsFabricUsable|createComm)[[:space:]]*\(' "$rev" -- '*.cpp' '*.h' '*.cu' '*.py' | head -180
echo "=== $rev NCCL env references ==="
git grep -n -E 'NCCL_(RUNTIME_CONNECT|GRAPH_REGISTER|NVLS_ENABLE)' "$rev" -- ':!*.lock'
done
echo "=== target affected dependency diff ==="
git diff "$base" "$target" -- cpp/tensorrt_llm/common/opUtils.cpp cpp/tensorrt_llm/common/ipUtils.cpp cpp/include/tensorrt_llm/runtime/ipcNvlsMemory.h cpp/tensorrt_llm/runtime/ipcNvlsMemory.cu
echo "=== three-way file preview ==="
git merge-file -p <(git show "$head:cpp/tensorrt_llm/common/opUtils.cpp") <(git show "$base:cpp/tensorrt_llm/common/opUtils.cpp") <(git show "$target:cpp/tensorrt_llm/common/opUtils.cpp")
echo "merge-file status=$?"Length of output: 102848 🏁 Script executed: #!/bin/bash
set -uo pipefail
head=20130077a0ba487470425a4af1a22b7b01ac9760
target=21d2f1c253c60da241c187f9489e882f29a31425
base=adea5c0720f266ef3ec8490b638498f491af682d
for f in cpp/tensorrt_llm/runtime/ipcNvlsMemory.cu cpp/tensorrt_llm/runtime/ncclCommunicator.cpp cpp/tensorrt_llm/thop/alltoallOp.cpp cpp/tests/unit_tests/runtime/ipcNvlsMemoryTest.cpp cpp/tests/unit_tests/multi_gpu/ncclUtilsTest.cpp; do
ast-grep outline "$f" --items all --match 'ipcNvls|MemHandle|createComm|Alltoall|SetUp|TEST|Nccl'
done
python - <<'PY'
import subprocess,tempfile,pathlib
h='20130077a0ba487470425a4af1a22b7b01ac9760'
t='21d2f1c253c60da241c187f9489e882f29a31425'
b='adea5c0720f266ef3ec8490b638498f491af682d'
f='cpp/tensorrt_llm/common/opUtils.cpp'
sources=[subprocess.check_output(['git','show',r+':'+f]) for r in (h,b,t)]
print('target opUtils identical to merge-base:', sources[1]==sources[2])
with tempfile.TemporaryDirectory(prefix='semantic-inspection-') as d:
paths=[]
for name,data in zip(('head','base','target'),sources):
p=pathlib.Path(d)/name;p.write_bytes(data);paths.append(str(p))
result=subprocess.run(['git','merge-file','-p',*paths],capture_output=True)
print('merge-file status:',result.returncode)
print('preview identical to head:',result.stdout==sources[0])
print(result.stderr.decode())
lines=result.stdout.decode().splitlines()
print('\n'.join(f'{i+1:4} {s}' for i,s in enumerate(lines) if 78<=i<=185))
PY
for rev in "$head" "$target"; do
echo "=== $rev fabric probe and communicator ==="
if [ "$rev" = "$head" ]; then
git show "$rev:cpp/tensorrt_llm/runtime/ipcNvlsMemory.cu" | nl -ba | sed -n '148,248p;475,548p'
else
git show "$rev:cpp/tensorrt_llm/runtime/ipcNvlsMemory.cu" | nl -ba | sed -n '224,324p;560,633p'
fi
git show "$rev:cpp/tensorrt_llm/runtime/ncclCommunicator.cpp" | nl -ba | sed -n '140,200p'
git show "$rev:cpp/tensorrt_llm/common/opUtils.h" | nl -ba | sed -n '173,189p'
echo "=== all getComm bindings including non-call references ==="
git grep -n -w getComm "$rev" -- ':!*.lock'
echo "=== alltoall caller ==="
git show "$rev:cpp/tensorrt_llm/thop/alltoallOp.cpp" | nl -ba | sed -n '18,105p'
echo "=== fabric tests ==="
git show "$rev:cpp/tests/unit_tests/runtime/ipcNvlsMemoryTest.cpp" | nl -ba | sed -n '24,130p'
echo "=== communicator tests ==="
git show "$rev:cpp/tests/unit_tests/multi_gpu/ncclUtilsTest.cpp" | nl -ba | sed -n '125,170p;265,300p;675,708p'
done
echo "=== relevant target caller edits ==="
git diff "$base" "$target" -- cpp/tensorrt_llm/thop/alltoallOp.cpp cpp/tensorrt_llm/common/attentionOp.h cpp/tensorrt_llm/thop/attentionOp.cpp tensorrt_llm/_torch/distributed/ops.py tensorrt_llm/_torch/moe/fused_moe/communication/nvlink_one_sided.py | head -350
echo "=== test path discovery at fixed revisions ==="
for rev in "$head" "$target"; do
git grep -n -E 'def test_configurable_moe_multi_gpu|alltoall|all_to_all' "$rev" -- 'tests/unittest/_torch/moe/*' | head -70
doneLength of output: 88754
SEMANTIC_REVIEW
Verdict: PASS. No concrete cross-branch interaction or PR-local compatibility defect was found in the inspected paths. Revision and integration verification All three supplied IDs resolve to commit objects. I compared both Inspected contracts and behavior
I also inspected Limitations This was source and Git inspection only. I did not execute project code or tests. I did not verify NCCL bootstrap timing or reproduce the reported intermittent failure. The review does not establish process-wide environment thread safety or certify unrelated target changes. SEMANTIC_RESULT request_id=919dd9c6-2f3b-4b0f-99e4-d15a44fa5a62 head=20130077a0ba487470425a4af1a22b7b01ac9760 target=21d2f1c253c60da241c187f9489e882f29a31425 merge_base=adea5c0720f266ef3ec8490b638498f491af682d verdict=PASS You are interacting with an AI system. |
|
PR_Github #75795 [ run ] completed with state
|
|
/bot run --disable-fail-fast |
|
PR_Github #75799 [ run ] triggered by Bot. Commit: |
|
PR_Github #75799 [ run ] completed with state |
Dev Engineer Review
getComm()now updates the NCCL environment before callinggetUniqueId(). This avoids concurrent environment updates during NCCL bootstrap. The updates preserve existing values. On non-Windows platforms, the code also disables NVLS whenipcNvlsFabricUsable()is false. No API changes are shown.QA Engineer Review
No test changes.
Per-File QA Perspective
cpp/tensorrt_llm/common/opUtils.cpp: Verify communicator initialization with unset and explicitly set NCCL variables. Verify that NVLS is disabled only when the fabric is unusable, and that communicator creation still succeeds through the fallback path.Description
Why. The TTP cases of
test_configurable_moe_multi_gpu(for exampleparallel=TTP-comm=NVLINK_ONE_SIDED-e8_k1_h512_i512-seq=8-dtype=torch.float16-backend=CUTLASS-quant=W8A16-routing=Renormalize) hang intermittently in CI on B200 and H100: 8 failures in 256 runs over the last 7 days, always on the first TTP case of a test group (NVBug 6801108). The group-root rank dies with SIGSEGV inside an NCCL thread, the other ranks wait forever in NCCL bootstrap, and pytest only stops at its timeout.getComm()incpp/tensorrt_llm/common/opUtils.cppcallssetenv()forNCCL_RUNTIME_CONNECT,NCCL_GRAPH_REGISTERand, when the NVLS fabric is not usable,NCCL_NVLS_ENABLEaftergetUniqueId(). On the group root,ncclGetUniqueId()starts NCCL'sbootstrapRootthread. Since NCCL 2.30.7 that thread callsgetenv("NCCL_SOCKET_MAGIC")the first time it initializes a socket:ncclSocketDefaultMagic()is new in 2.30.7 and is now the defaultmagicargument ofncclSocketInit(), whereas 2.30.4 used the constantNCCL_SOCKET_MAGIC.setenv()is not thread-safe, because adding a variable can reallocateenvironwhile the bootstrap thread is scanning it. The CI image moved from NCCL 2.30.4 to 2.30.7 in #18126, which is when the hang started on main.Native stack of the crashing thread, captured on the group root of a hung CI-equivalent run:
What changed. The environment updates in
getComm()now run beforegetUniqueId(), so nosetenv()from this function can overlap with the NCCL bootstrap thread it starts. This is the orderNcclCommunicator::createComm()already uses. The same variables are set with the same no-overwrite semantics, so there is no other behavior change.Test Coverage
No new test: the failure is a timing-dependent race. The existing test exercises this path and was used to validate the change.
unittest/_torch/moe/test_moe_module.py::test_configurable_moe_multi_gpu -k "CUTLASS and W8A16"(l0_dgx_b200.yml) was run in a loop through the CI harness (test_unittests_v2) on a B200x4 node, with the same image, install steps, and pytest command asDGX_B200-4_GPUs-PyTorch-1:33193c2b, without this change)setenv()calls no-ops33193c2b+ this change, rebuilt with the CI cmake options (SM100 only)At the CI failure rate (about 3.1% per run), 125 consecutive clean rounds would occur with probability of about 2% without the fix.
A standalone reproducer (no TensorRT-LLM, no GPU) starts a fresh process that calls
ncclGetUniqueId()and then adds 256 environment variables withsetenv():setenv()relative toncclGetUniqueId()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)
If PR introduces API changes, an appropriate PR label is added - either
api-compatibleorapi-breaking. Forapi-breaking, includeBREAKINGin the PR title.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.