[None][fix] Fall back to NCCL when CUDA IPC handles cannot be exchanged - #17034
[None][fix] Fall back to NCCL when CUDA IPC handles cannot be exchanged#17034pjdurden wants to merge 3 commits into
Conversation
|
Note Reviews pausedIt looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the Use the following commands to manage reviews:
Use the checkboxes below for quick actions:
WalkthroughCUDA IPC initialization now coordinates failures across tensor-parallel ranks, cleans up partial resources, and reports unavailable IPC. Allreduce initialization selects NCCL or raises an error based on IPC availability. Tests cover success, failure, fallback, and cleanup paths. ChangesCUDA IPC fallback
Estimated code review effort: 4 (Complex) | ~45 minutes Suggested reviewers: Sequence Diagram(s)sequenceDiagram
participant AllReduce
participant IpcMemory
participant TPCollective
participant NCCL
AllReduce->>IpcMemory: initialize IPC workspace
IpcMemory->>TPCollective: exchange IPC failure status
TPCollective-->>IpcMemory: return TP-wide status
IpcMemory-->>AllReduce: return workspace and failure flag
AllReduce->>NCCL: select NCCL when IPC failed
🚥 Pre-merge checks | ✅ 4 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (4 passed)
✨ Finishing Touches🧪 Generate unit tests (beta)
Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 1
🧹 Nitpick comments (2)
tests/unittest/_torch/distributed/test_ipc_memory_fallback.py (1)
74-157: 📐 Maintainability & Code Quality | 🔵 Trivial | 🏗️ Heavy liftExpand coverage for the remaining fallback contracts.
Coverage summary — insufficient. The three added tests cover successful IPC setup, local import failure, and peer-reported failure. CI/QA test-list membership cannot be verified because no
tests/integration/test_lists/test-db/ortests/integration/test_lists/qa/file was provided.
- Make
FakeCudart.cudaIpcGetMemHandle()configurable and test its new dummy-handle/cleanup path.- In
test_local_ipc_open_failure_falls_back_instead_of_raising, useFakeDist(TP_SIZE)so a local failure alone drives agreement;peer_ipc_ok=Falsecurrently masks failures to propagate the local error.- Add mocked coverage that unavailable IPC downgrades
AllReduceto NCCL withworkspace is None, and thatlamport_initialize()is skipped.As per path instructions, test changes require coverage and test-list status review.
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tests/unittest/_torch/distributed/test_ipc_memory_fallback.py` around lines 74 - 157, Expand the IPC fallback tests around FakeCudart.cudaIpcGetMemHandle, making its result configurable and covering the dummy-handle and cleanup behavior. Update test_local_ipc_open_failure_falls_back_instead_of_raising to use FakeDist(TP_SIZE), then add mocked coverage verifying unavailable IPC downgrades AllReduce to NCCL with workspace=None and skips lamport_initialize(); also review coverage and test-list status for these test changes.Source: Path instructions
tensorrt_llm/_ipc_utils.py (1)
18-18: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winUse built-in generics in
open_ipc_memoryReplace
Optional[Tuple[List[int], int]]withtuple[list[int], int] | None.Listis still needed by the other annotations in this file, so thetypingimport can’t be dropped yet.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tensorrt_llm/_ipc_utils.py` at line 18, Update the return annotation of open_ipc_memory to use the built-in generic form tuple[list[int], int] | None instead of Optional[Tuple[List[int], int]]. Retain the existing typing imports because List remains required by other annotations in the file.Sources: Coding guidelines, Learnings
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Inline comments:
In `@tensorrt_llm/_ipc_utils.py`:
- Around line 161-163: Update the CUDA IPC warning in the logger.warning call to
use the f-string conversion flag for ipc_error instead of explicitly calling
repr(), resolving Ruff RUF010 while preserving the existing message.
---
Nitpick comments:
In `@tensorrt_llm/_ipc_utils.py`:
- Line 18: Update the return annotation of open_ipc_memory to use the built-in
generic form tuple[list[int], int] | None instead of Optional[Tuple[List[int],
int]]. Retain the existing typing imports because List remains required by other
annotations in the file.
In `@tests/unittest/_torch/distributed/test_ipc_memory_fallback.py`:
- Around line 74-157: Expand the IPC fallback tests around
FakeCudart.cudaIpcGetMemHandle, making its result configurable and covering the
dummy-handle and cleanup behavior. Update
test_local_ipc_open_failure_falls_back_instead_of_raising to use
FakeDist(TP_SIZE), then add mocked coverage verifying unavailable IPC downgrades
AllReduce to NCCL with workspace=None and skips lamport_initialize(); also
review coverage and test-list status for these test changes.
🪄 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: Enterprise
Run ID: ab9b4213-0365-4477-a4b7-557a34646fe2
📒 Files selected for processing (4)
tensorrt_llm/_ipc_utils.pytensorrt_llm/_torch/distributed/allreduce_helper.pytensorrt_llm/_torch/distributed/ops.pytests/unittest/_torch/distributed/test_ipc_memory_fallback.py
There was a problem hiding this comment.
Pull request overview
This PR addresses a Tensor Parallel initialization failure where CUDA IPC handle import/export can fail (e.g., cudaIpcOpenMemHandle returning cudaErrorInvalidDevice) even when cudaDeviceCanAccessPeer reports P2P capability, by degrading to the already-supported NCCL path instead of aborting engine initialization.
Changes:
- Make
IpcMemory.open_ipc_memory()returnNoneon CUDA IPC handle exchange failures (with TP-wide agreement + cleanup), and haveIpcMemory.__init__disable IPC (null pointers) instead of raising. - Gate Lamport buffer initialization on
lamport_buffers.open_ipcrather than the optimistic P2P check. - In
AllReduce.__init__, detect workspaces that have no usable IPC buffers and downgrade the strategy to NCCL; add a focused unit test module with stubbed CUDA runtime / distributed collectives.
Reviewed changes
Copilot reviewed 4 out of 4 changed files in this pull request and generated 2 comments.
| File | Description |
|---|---|
tensorrt_llm/_ipc_utils.py |
Make CUDA IPC failures non-fatal by returning None and cleaning up; disables IPC state in IpcMemory on failure. |
tensorrt_llm/_torch/distributed/allreduce_helper.py |
Prevent lamport_initialize() from running when IPC buffers are not actually available. |
tensorrt_llm/_torch/distributed/ops.py |
Add IPC-workspace detection and force NCCL fallback when IPC buffers are missing; reorder LOWPRECISION allocation to avoid null-workspace usage. |
tests/unittest/_torch/distributed/test_ipc_memory_fallback.py |
Add stubbed unit tests covering IPC-success, local IPC-open failure, and peer IPC failure agreement behavior. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
| self.workspace = get_allreduce_workspace(self.mapping) | ||
| # Every custom all-reduce kernel reads the peers' IPC buffers. When | ||
| # CUDA IPC is unavailable those pointers are null, so NCCL is the | ||
| # only strategy that can run. See allreduce_workspace_has_ipc. | ||
| if not allreduce_workspace_has_ipc(self.mapping): |
| if not all(dist.tp_allgather(ipc_error is None)): | ||
| if ipc_error is not None: | ||
| logger.warning( | ||
| f"CUDA IPC is not usable on this system: {repr(ipc_error)}. " | ||
| "Custom all-reduce kernels are disabled, falling back to NCCL." |
|
@BowenFu @schetlur-nv @JunyiXu-nv @allisonlim-nv could one of you take a look when you get a chance? The reporter of #16899 applied the patch on 8x RTX 5090 (PCIe, no NVLink), a different platform from the RTX 6000D in the original report: Before: crash at engine init, cudaErrorInvalidDevice: 101 from cudaIpcOpenMemHandle, all 8 ranks. So the failure is not specific to the RTX 6000D, it affects PCIe-only parts generally. Also pushed one more commit for Ruff nit. |
|
/bot run |
brnguyen2
left a comment
brnguyen2
left a comment
There was a problem hiding this comment.
The _ipc_utils.py fix is the right one and the collective-agreement + cleanup details are handled carefully. Two things before merge.
The ops.py downgrade is broader than the bug. allreduce_workspace_has_ipc() is False whenever can_access_peer() is False — which includes the entirely normal inter-node TP case (can_access_peer returns False on a node-rank mismatch, and IpcMemory.__init__ already forces open_ipc=False when tp_size > gpus_per_node). See the inline comment at [ops.py:785](https://github.com/NVIDIA/TensorRT-LLM/pull/17034/files#diff-8343f10ef2d25d7245cbff50bea93ad47fd865db34a1e3c1d116ae57f7ab7ac1R785); as written this disables MNNVL on multi-node NVLink systems. The _ipc_utils + allreduce_helper changes alone already fix #16899; if you keep the ops.py guard, it should trigger only on the new failure mode (P2P reported, IPC unusable) and run after the MNNVL block.
Test coverage stops at _ipc_utils. The three new tests are good, but nothing covers the two behavioral changes that can break working hardware: the AllReduce.__init__ strategy downgrade and the lamport_initialize gating. A test asserting that an MNNVL-eligible mapping still gets mnnvl_allreduce set would have caught the issue above. tests/unittest/_torch/distributed is enrolled as a directory in l0_dgx_h100.yml, so the new file will run — worth confirming the cudart patching is safe alongside the real-GPU tests in that stage.
Description nits: the last paragraph is cut off mid-sentence ("trading a "), and the claim that FP8 "worked because it fits in fewer bytes, not because it dodges IPC" contradicts the following clause about not hitting the init path. Also, MoEAllReduce and MiniMaxAllReduceRMS still take the workspace unconditionally, so the "defense in depth" argument doesn't cover them — worth saying so explicitly rather than implying full coverage.
| # Every custom all-reduce kernel reads the peers' IPC buffers. When | ||
| # CUDA IPC is unavailable those pointers are null, so NCCL is the | ||
| # only strategy that can run. See allreduce_workspace_has_ipc. | ||
| if not allreduce_workspace_has_ipc(self.mapping): |
There was a problem hiding this comment.
This fires far more often than the bug it targets. allreduce_workspace_has_ipc() is False for any workspace built with is_p2p_supported=False, and can_access_peer() returns False on a plain node-rank mismatch ([_ipc_utils.py:46](https://github.com/NVIDIA/TensorRT-LLM/pull/17034/files#diff-58a9acafed7455cb039511131ef7c8fb87d8d39200b75e733c93914b25e26eecR46)) — plus IpcMemory.__init__ forces open_ipc=False when tp_size > gpus_per_node. So on any multi-node TP config this rewrites self.strategy to NCCL, and the MNNVL block right below ([ops.py:795](https://github.com/NVIDIA/TensorRT-LLM/pull/17034/files#diff-8343f10ef2d25d7245cbff50bea93ad47fd865db34a1e3c1d116ae57f7ab7ac1R795)) is keyed on self.strategy in (AUTO, MNNVL) — so mnnvl_allreduce is never constructed. That's a silent perf regression on GB200-class multi-node NVLink, which is exactly the topology MNNVL exists for.
Two fixes, either works: (a) only downgrade when the new failure mode occurred — i.e. can_access_peer(mapping) was True but open_ipc came back False — so systems that never had P2P keep today's behavior; or (b) move this block below the MNNVL init and skip it when self.mnnvl_allreduce is not None. (a) is narrower and matches the PR's stated scope.
There was a problem hiding this comment.
Went with (a). IpcMemory tracks ipc_failed separately from open_ipc now, set only when IPC was requested and the group fits in one node but the exchange still failed. can_access_peer False, node-rank mismatch, tp_size > gpus_per_node all leave it False, so multi-node TP does not touch self.strategy and the MNNVL block runs like before.
Added test_mnnvl_kept_without_p2p for it.
| # CUDA IPC is unavailable those pointers are null, so NCCL is the | ||
| # only strategy that can run. See allreduce_workspace_has_ipc. | ||
| if not allreduce_workspace_has_ipc(self.mapping): | ||
| logger.warning( |
There was a problem hiding this comment.
AllReduce is instantiated per decoder layer (often several per layer), so on an affected system this prints a warning ~100+ times per rank. Use logger.warning_once(..., key=...), or better, emit the message once from open_ipc_memory where the condition is actually detected.
There was a problem hiding this comment.
Moved it. warning_once with a key, out of open_ipc_memory where the condition is detected, so one line per process instead of one per AllReduce. Left a logger.debug at the strategy site that names what got downgraded.
| ipc_error = None if error == cudart.cudaError_t.cudaSuccess else error | ||
| if ipc_error is not None: | ||
| # Exchange a dummy handle to keep the collective below symmetric. | ||
| local_handle = cudart.cudaIpcMemHandle_t() |
There was a problem hiding this comment.
Exchanging an all-zero cudaIpcMemHandle_t means every other rank calls cudaIpcOpenMemHandle on garbage. It probably just returns an error you catch below, but you're feeding an uninitialized handle to the driver to preserve collective symmetry when a simpler option exists: allgather the cudaIpcGetMemHandle success flag first and skip the handle exchange + open loop entirely if any rank failed. Same number of collectives (you need a second agreement allgather for open failures either way), no garbage handles.
There was a problem hiding this comment.
Dropped the dummy handle. A rank that cannot export sends None in the handle allgather, and everyone bails before the open loop if any entry is None. That collective carries the payload and the export agreement together, so still two collectives, second one covers open failures.
| peer_ptrs.append(ptr) | ||
| opened_ptrs.append(ptr) | ||
|
|
||
| if not all(dist.tp_allgather(ipc_error is None)): |
There was a problem hiding this comment.
Only ranks that failed locally log; ranks that succeeded and are falling back because a peer failed degrade silently. Since the group decision is what matters operationally, log unconditionally in this branch (with the peer-driven case worded differently), so a user reading rank 0's log can tell why custom all-reduce is off.
There was a problem hiding this comment.
Both branches log now, through a shared disable_ipc(). Message is either failed here with or failed on a peer rank of the TP group.
| still cannot be imported. In that case the workspace pointers are null and only | ||
| NCCL can be used. get_allreduce_workspace must have been called first. | ||
| """ | ||
| allreduce_workspaces = getattr(_thread_local, |
There was a problem hiding this comment.
getattr without a default raises AttributeError if get_allreduce_workspace hasn't run for this pp_rank — the docstring notes the precondition, but an unguarded getattr/[mapping] turns a misuse into an opaque traceback. Cheap fix: take the workspace tuple as an argument, or have get_allreduce_workspace return both the tensor and the flag so the invariant is structural rather than documented.
There was a problem hiding this comment.
get_allreduce_workspace returns (workspace, ipc_failed) and computes the flag when it fills the thread-local, so no second reader and no bare getattr. MoEAllReduce and MiniMaxAllReduceRMS go through _require_ipc_workspace() which raises, they have no NCCL path.
Addresses review feedback on NVIDIA#17034. allreduce_workspace_has_ipc() was False for any workspace built without P2P, which includes ordinary inter-node TP, so it rewrote the strategy to NCCL before the MNNVL block and mnnvl_allreduce was never constructed on multi-node NVLink systems. IpcMemory now records ipc_failed separately from open_ipc, and only that failure mode, P2P reported but handles unusable, downgrades the strategy. Configurations that never had P2P keep the behaviour they had. get_allreduce_workspace returns the flag alongside the workspace, so a caller cannot read the thread-local state before it has been populated. The fallback is reported once, from open_ipc_memory where it is detected, and on every rank of the group rather than only on the ones that failed locally. AllReduce is constructed per decoder layer, so it no longer warns per instance. A rank that cannot export its handle contributes None to the allgather instead of an uninitialized cudaIpcMemHandle_t, so no rank imports a garbage handle. The number of collectives is unchanged. MoEAllReduce and MiniMaxAllReduceRMS have no NCCL path, so they now raise an actionable error instead of letting the kernel dereference null peer pointers. Tests cover the strategy downgrade, MNNVL still being selected when the workspace has no IPC buffers by design, and the lamport_initialize gating. Signed-off-by: pjdurden <prajjwalchittori1@gmail.com>
|
Note GitHub couldn't provide a complete incremental comparison for this pull request, so CodeRabbit is performing a full review instead. This review may take a little longer. |
pjdurden
changed the title
|
All five inline points addressed. Pushed as a separate commit on top of the original fix so the delta reads on its own, then re-signed the branch to clear DCO, it was failing on the first commit. ops.py downgrade: went with (a). IpcMemory tracks ipc_failed separately from open_ipc, set only when IPC was requested and possible but the exchange failed. can_access_peer returning False leaves it False, so inter-node TP keeps current behaviour and still reaches the MNNVL block. get_allreduce_workspace returns (workspace, ipc_failed), which also kills the bare getattr at ops.py:62. Tests are 11 now. New: test_mnnvl_kept_without_p2p (the guard for the above), test_downgrade_to_nccl, test_lamport_skipped which drives the real allocate_allreduce_fusion_workspace and asserts lamport_initialize is not called, test_moe_allreduce_raises, test_peer_export_fails, plus two for ipc_failed staying False. On the cudart patching in that stage: every patch is unittest.mock.patch scoped to a with block on module attributes, nothing survives the test, and the fallback frees its fake pointers through the fake runtime before the patch unwinds. test_safe_mpi_comm.py in the same dir already patches this way. Only test_lamport_skipped touches real CUDA, it is gated on torch.cuda.is_available() because the helper allocates device tensors. MoEAllReduce and MiniMaxAllReduceRMS: you are right that defense in depth overclaimed. Neither has a non-IPC path, so they raise explicitly when ipc_failed is set. Same init-time failure they had before this PR, just with a message instead of a null deref in the kernel. Called out in the description now. Description rewritten. Truncated paragraph gone, FP8 claim dropped, I have no evidence for the mechanism there, only that those runs did not hit this path. Title follows [None][fix] now. Same caveat as before: no GPU or torch here so pytest did not run. The _ipc_utils logic is covered by a standalone harness that loads the real edited module with the CUDA runtime and collectives stubbed, six scenarios, all pass. Section 6 of the description lists them. Rest needs CI. DCO is green. Can you kick off another /bot run? |
There was a problem hiding this comment.
Actionable comments posted: 2
🧹 Nitpick comments (1)
tests/unittest/_torch/distributed/test_ipc_memory_fallback.py (1)
271-302: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winTest coverage summary: insufficient.
Added tests cover IPC success, local and peer IPC failures, intentional IPC absence, MNNVL preservation, NCCL fallback, MoE failure, and Lamport initialization.
Add a direct
MiniMaxAllReduceRMSconstruction test withipc_failed=True. The production change now routes this class through_require_ipc_workspace(), but no test verifies that integration.Test-list membership is not applicable to this review because no
tests/integration/test_lists/file was supplied. Runpytest tests/unittest/after the update.As per path instructions, test-code changes must include a coverage verdict.
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tests/unittest/_torch/distributed/test_ipc_memory_fallback.py` around lines 271 - 302, The tests lack direct coverage for MiniMaxAllReduceRMS when IPC is unavailable. Add a unit test that constructs MiniMaxAllReduceRMS with ipc_failed=True and verifies the production path raises the expected requires-CUDA-IPC error through _require_ipc_workspace(); then run pytest tests/unittest/ and include a coverage verdict for the added test.Sources: Coding guidelines, Path instructions
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Inline comments:
In `@tensorrt_llm/_torch/distributed/ops.py`:
- Around line 69-81: Update _require_ipc_workspace to distinguish IPC usability
from ipc_failed, and reject workspaces when IPC is intentionally unavailable as
well as when initialization fails; do not return a workspace containing null
peer pointers to MoEAllReduce or MiniMaxAllReduceRMS. Preserve the upstream
contract that intentional IPC absence leaves ipc_failed false, and add
inter-node coverage for both IPC-only fused operations.
In `@tests/unittest/_torch/distributed/test_ipc_memory_fallback.py`:
- Around line 140-143: Update _run so the IpcMemory created with open_ipc=True
is explicitly cleaned up while cudart_patch and dist_patch remain active, then
set its open_ipc attribute to False before returning it to prevent
IpcMemory.__del__ from performing a second cleanup.
---
Nitpick comments:
In `@tests/unittest/_torch/distributed/test_ipc_memory_fallback.py`:
- Around line 271-302: The tests lack direct coverage for MiniMaxAllReduceRMS
when IPC is unavailable. Add a unit test that constructs MiniMaxAllReduceRMS
with ipc_failed=True and verifies the production path raises the expected
requires-CUDA-IPC error through _require_ipc_workspace(); then run pytest
tests/unittest/ and include a coverage verdict for the added test.
🪄 Autofix
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: Enterprise
Run ID: f1588ceb-fd00-410f-8baf-31aafa0793da
📒 Files selected for processing (4)
tensorrt_llm/_ipc_utils.pytensorrt_llm/_torch/distributed/allreduce_helper.pytensorrt_llm/_torch/distributed/ops.pytests/unittest/_torch/distributed/test_ipc_memory_fallback.py
🚧 Files skipped from review as they are similar to previous changes (2)
- tensorrt_llm/_torch/distributed/allreduce_helper.py
- tensorrt_llm/_ipc_utils.py
| def _require_ipc_workspace(mapping: Mapping, op_name: str) -> torch.LongTensor: | ||
| """Workspace for fused ops that reinterpret it as `void**` and have no NCCL path. | ||
|
|
||
| AllReduce degrades to NCCL when CUDA IPC turns out to be unusable, but these | ||
| kernels cannot: they would dereference the null peer pointers. Fail here with an | ||
| actionable message rather than in the kernel with an illegal memory access. | ||
| """ | ||
| workspace, ipc_failed = get_allreduce_workspace(mapping) | ||
| if ipc_failed: | ||
| raise RuntimeError( | ||
| f"{op_name} requires CUDA IPC, which is unavailable on this system, " | ||
| "and it has no NCCL fallback.") | ||
| return workspace |
There was a problem hiding this comment.
🩺 Stability & Availability | 🟠 Major | 🏗️ Heavy lift
Reject workspaces with no usable IPC buffers.
ipc_failed is false when IPC was intentionally unavailable, such as inter-node TP. In that state, IpcMemory still serializes null peer pointers. Lines 76-81 return this workspace to MoEAllReduce and MiniMaxAllReduceRMS, and their fused kernels can dereference the null pointers.
Track IPC usability separately from IPC failure. Make _require_ipc_workspace() reject both states. Add inter-node tests for both IPC-only fused operations.
Based on the supplied upstream contract, intentionally absent IPC keeps ipc_failed=False and leaves peer pointers null.
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
In `@tensorrt_llm/_torch/distributed/ops.py` around lines 69 - 81, Update
_require_ipc_workspace to distinguish IPC usability from ipc_failed, and reject
workspaces when IPC is intentionally unavailable as well as when initialization
fails; do not return a workspace containing null peer pointers to MoEAllReduce
or MiniMaxAllReduceRMS. Preserve the upstream contract that intentional IPC
absence leaves ipc_failed false, and add inter-node coverage for both IPC-only
fused operations.
| def _run(mapping, cudart, dist, open_ipc=True): | ||
| cudart_patch, dist_patch = _patched(cudart, dist) | ||
| with cudart_patch, dist_patch: | ||
| return IpcMemory(mapping, 1 << 20, open_ipc) |
There was a problem hiding this comment.
🩺 Stability & Availability | 🟡 Minor | ⚡ Quick win
Clean up the fake IPC allocation before the patches exit.
_run() returns an IpcMemory with open_ipc=True after it restores _ipc_utils.cudart. When the successful test releases that object, IpcMemory.__del__() can call the real CUDA runtime with fake pointers.
Keep the patches active through explicit IPC cleanup. Then set open_ipc to False to prevent a second cleanup in __del__.
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
In `@tests/unittest/_torch/distributed/test_ipc_memory_fallback.py` around lines
140 - 143, Update _run so the IpcMemory created with open_ipc=True is explicitly
cleaned up while cudart_patch and dist_patch remain active, then set its
open_ipc attribute to False before returning it to prevent IpcMemory.__del__
from performing a second cleanup.
Addresses review feedback on NVIDIA#17034. allreduce_workspace_has_ipc() was False for any workspace built without P2P, which includes ordinary inter-node TP, so it rewrote the strategy to NCCL before the MNNVL block and mnnvl_allreduce was never constructed on multi-node NVLink systems. IpcMemory now records ipc_failed separately from open_ipc, and only that failure mode, P2P reported but handles unusable, downgrades the strategy. Configurations that never had P2P keep the behaviour they had. get_allreduce_workspace returns the flag alongside the workspace, so a caller cannot read the thread-local state before it has been populated. The fallback is reported once, from open_ipc_memory where it is detected, and on every rank of the group rather than only on the ones that failed locally. AllReduce is constructed per decoder layer, so it no longer warns per instance. A rank that cannot export its handle contributes None to the allgather instead of an uninitialized cudaIpcMemHandle_t, so no rank imports a garbage handle. The number of collectives is unchanged. MoEAllReduce and MiniMaxAllReduceRMS have no NCCL path, so they now raise an actionable error instead of letting the kernel dereference null peer pointers. Tests cover the strategy downgrade, MNNVL still being selected when the workspace has no IPC buffers by design, and the lamport_initialize gating. Signed-off-by: pjdurden <prajjwalchittori1@gmail.com>
|
Note GitHub couldn't provide a complete incremental comparison for this pull request, so CodeRabbit is performing a full review instead. This review may take a little longer. |
|
Note GitHub couldn't provide a complete incremental comparison for this pull request, so CodeRabbit is performing a full review instead. This review may take a little longer. |
Addresses review feedback on NVIDIA#17034. allreduce_workspace_has_ipc() was False for any workspace built without P2P, which includes ordinary inter-node TP, so it rewrote the strategy to NCCL before the MNNVL block and mnnvl_allreduce was never constructed on multi-node NVLink systems. IpcMemory now records ipc_failed separately from open_ipc, and only that failure mode, P2P reported but handles unusable, downgrades the strategy. Configurations that never had P2P keep the behaviour they had. get_allreduce_workspace returns the flag alongside the workspace, so a caller cannot read the thread-local state before it has been populated. The fallback is reported once, from open_ipc_memory where it is detected, and on every rank of the group rather than only on the ones that failed locally. AllReduce is constructed per decoder layer, so it no longer warns per instance. A rank that cannot export its handle contributes None to the allgather instead of an uninitialized cudaIpcMemHandle_t, so no rank imports a garbage handle. The number of collectives is unchanged. MoEAllReduce and MiniMaxAllReduceRMS have no NCCL path, so they now raise an actionable error instead of letting the kernel dereference null peer pointers. Tests cover the strategy downgrade, MNNVL still being selected when the workspace has no IPC buffers by design, and the lamport_initialize gating. Signed-off-by: pjdurden <prajjwalchittori1@gmail.com>
|
Rebased onto main. It was 162 behind, now 0, and it was a clean rebase, none of One thing that changed under us and is worth flagging, since your review note assumed otherwise: #16498 moved
On the Separately, checking the MNNVL interaction more closely: |
There was a problem hiding this comment.
Actionable comments posted: 2
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Inline comments:
In `@tests/unittest/_torch/distributed/test_ipc_memory_fallback.py`:
- Around line 92-98: Update FakeCudart.cudaMalloc and FakeCudart.cudaMemset to
support configurable CUDA failures, then add tests covering each RuntimeError
path and verifying errors propagate without falling back from IPC to NCCL. Keep
the existing success behavior unchanged and exercise the failure configuration
through the relevant distributed IPC test flows.
- Around line 46-51: Update _has_allreduce_op to catch only AttributeError when
accessing torch.ops.trtllm.allreduce, allowing unrelated setup failures to
propagate while preserving the false result for a missing operator.
🪄 Autofix
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: Enterprise
Run ID: 39b3239f-8583-45ba-9fca-ce8317e5b2d7
📒 Files selected for processing (4)
tensorrt_llm/_ipc_utils.pytensorrt_llm/_torch/distributed/allreduce_helper.pytensorrt_llm/_torch/distributed/ops.pytests/unittest/_torch/distributed/test_ipc_memory_fallback.py
🚧 Files skipped from review as they are similar to previous changes (3)
- tensorrt_llm/_torch/distributed/allreduce_helper.py
- tensorrt_llm/_ipc_utils.py
- tensorrt_llm/_torch/distributed/ops.py
| def _has_allreduce_op() -> bool: | ||
| """Whether the trtllm custom ops are registered in this build.""" | ||
| try: | ||
| return torch.ops.trtllm.allreduce is not None | ||
| except Exception: | ||
| return False |
There was a problem hiding this comment.
🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win
🧩 Analysis chain
🌐 Web query:
In the supported PyTorch versions for TensorRT-LLM, which exception does torch.ops..<missing_operator> raise when the operator is not registered?
💡 Result:
When using PyTorch's torch.ops to access an operator that has not been registered, the behavior depends on the specific API implementation used to retrieve the operator, but it typically raises a ValueError or a RuntimeError [1][2][3]. According to the PyTorch source code for operator resolution (e.g., torch/_custom_op/impl.py), the function used to retrieve an operator raises a ValueError with the message: "Could not find the operator . Please make sure you have already registered the operator and (if registered from C++) loaded it via torch.ops.load_library." [3]. This specific error is frequently encountered when an operator is missing or not properly loaded [1]. In other contexts—such as when the dispatcher attempts to resolve a call to an operator that does not exist—PyTorch may raise a RuntimeError, often phrased as "RuntimeError: operator does not exist" or "RuntimeError: No such operator " [4][2][5]. While TensorRT-LLM relies on PyTorch's operator infrastructure, these exceptions are standard PyTorch behaviors triggered when the underlying dispatcher or operator registry cannot resolve the requested operator [3][6]. If you encounter this while using TensorRT-LLM, it usually indicates that the necessary library (e.g., a custom C++ extension or kernel library) has not been loaded or that there is a version mismatch between the environment and the compiled extensions [1][4][5].
Citations:
- 1: "ValueError: Could not find the operator torchvision::nms" upon importing torchvision pytorch/vision#8289
- 2: RuntimeError: operator torchvision::nms does not exist pytorch/pytorch#157648
- 3: https://github.com/pytorch/pytorch/blob/ec673ecd/torch/_custom_op/impl.py
- 4: Custom operator registration doesn't work when libtorch.so changes pytorch/pytorch#165151
- 5: API torch.ops.image.read_file reports RuntimeError - No such operator image::read_file pytorch/pytorch#65473
- 6: pytorch backend run error with fp8 hf model #2825
🏁 Script executed:
#!/bin/bash
set -eu
printf '%s\n' '--- target file outline ---'
ast-grep outline tests/unittest/_torch/distributed/test_ipc_memory_fallback.py --match '$_' --view compact 2>/dev/null || true
printf '%s\n' '--- helper and related usage ---'
rg -n -C 5 '_has_allreduce_op|allreduce|pytest\.skip|except ' tests/unittest/_torch/distributed/test_ipc_memory_fallback.py
printf '%s\n' '--- installed PyTorch behavior probe ---'
python3 - <<'PY'
try:
import torch
except Exception as exc:
print(f"torch import unavailable: {type(exc).__name__}: {exc}")
else:
print(f"torch version: {torch.__version__}")
for namespace, operator in (("trtllm", "allreduce"), ("definitely_missing_namespace", "definitely_missing_operator")):
try:
value = getattr(getattr(torch.ops, namespace), operator)
except BaseException as exc:
print(f"{namespace}.{operator}: {type(exc).__name__}: {exc}")
else:
print(f"{namespace}.{operator}: returned {value!r}")
PYRepository: NVIDIA/TensorRT-LLM
Length of output: 4683
🌐 Web query:
site:github.com/pytorch/pytorch "_OpNamespace" "__getattr__" "has no attribute"
💡 Result:
The error AttributeError: '_OpNamespace' '[namespace]' object has no attribute '[op_name]' is a PyTorch error that occurs when you attempt to access an operator within a namespace (e.g., torch.ops.aten, torch.ops._C) that has not been registered or is otherwise inaccessible [1][2]. This error originates from the _OpNamespace class in PyTorch, which dynamically handles operator lookups [1]. When you call an attribute on a namespace, the getattr method attempts to retrieve the specified operator from the PyTorch C++ operator registry [1][2]. If the operator cannot be found, an AttributeError is raised [1][2]. Common causes for this error include: 1. Unregistered or Missing Custom Operators: If you are using a custom operator (common in projects like vLLM), the underlying C++ extension that registers that operator may not have been correctly built, loaded, or installed [3][4]. 2. Version Mismatches: The operator may have been renamed, moved, or removed in the version of PyTorch you are using, or you may be calling an operator that is not present in your specific PyTorch build [5][6][7]. 3. Loading Issues: For C++ extensions, the library might not be fully loaded before you attempt to access the operator. Ensure that any custom libraries or extensions are initialized before calling their functions [8][9]. 4. Incorrect Namespace or Name: Ensure you are using the correct namespace and operator name. In some cases, extensions might use a different namespace than expected or require specific macro expansion during registration [8][9]. To troubleshoot this: - Verify that the extension or library containing the operator is correctly installed and imported. - Check the PyTorch version compatibility, especially if the code was written for a different version. - If using third-party packages (like vLLM), verify that the package is compatible with your PyTorch environment [3][4]. - If you are developing the custom operator, verify your C++ registration code (e.g., TORCH_LIBRARY) to ensure the namespace is correctly defined [8].
Citations:
- 1: https://github.com/pytorch/pytorch/blob/main/torch/_ops.py
- 2: https://github.com/pytorch/pytorch/blob/b2311192e6c4745aac3fdd774ac9d56a36b396d4/torch/_ops.py
- 3: torch._dynamo.exc.InternalTorchDynamoError: AttributeError: '_OpNamespace' '_C' object has no attribute 'rotary_embedding' pytorch/pytorch#161735
- 4: AttributeError: '_OpNamespace' '_C' object has no attribute 'silu_and_mul' pytorch/pytorch#146785
- 5: AttributeError: '_OpNamespace' 'aten' object has no attribute '_int_mm' pytorch/pytorch#98176
- 6: AttributeError: '_OpNamespace' 'aten' object has no attribute 'momentum' pytorch/pytorch#145274
- 7: torch=2.1, import torch; fails AttributeError: '_OpNamespace' 'aten' object has no attribute 'sym_constrain_range_for_size' pytorch/pytorch#127388
- 8: TORCH_LIBRARY breaks when passing (unexpanded) macro as namespace argument pytorch/pytorch#127615
- 9: Custom Symbolic Function Registration for ONNX Error pytorch/pytorch#105160
🏁 Script executed:
#!/bin/bash
set -eu
printf '%s\n' '--- PyTorch version constraints in the repository ---'
rg -n -i -C 2 'torch([<>=!~]|[-_ ]version)|pytorch([<>=!~]|[-_ ]version)' \
pyproject.toml setup.py setup.cfg requirements* docker .github 2>/dev/null | head -200 || true
printf '%s\n' '--- local references to the operator helper or skip marker ---'
rg -n -C 3 '_has_allreduce_op|requires_allreduce_op|allreduce custom op' .Repository: NVIDIA/TensorRT-LLM
Length of output: 7845
🏁 Script executed:
#!/bin/bash
set -eu
for ref in v2.11.0 main; do
url="https://raw.githubusercontent.com/pytorch/pytorch/${ref}/torch/_ops.py"
printf '%s\n' "--- ${ref} ---"
curl -fsSL "$url" |
grep -n -A35 -B5 'class _OpNamespace' |
head -80
doneRepository: NVIDIA/TensorRT-LLM
Length of output: 4248
🏁 Script executed:
#!/bin/bash
set -eu
for ref in v2.11.0 main; do
printf '%s\n' "--- ${ref} missing-operator branch ---"
curl -fsSL "https://raw.githubusercontent.com/pytorch/pytorch/${ref}/torch/_ops.py" |
sed -n '/def __getattr__(self, op_name: str)/,+32p'
doneRepository: NVIDIA/TensorRT-LLM
Length of output: 3278
🏁 Script executed:
#!/bin/bash
set -eu
printf '%s\n' '--- changed files ---'
git diff --name-only
printf '%s\n' '--- target test functions ---'
rg -n '^def test_|^async def test_' tests/unittest/_torch/distributed/test_ipc_memory_fallback.py
printf '%s\n' '--- matching test-list entries ---'
rg -n -i 'ipc_memory_fallback|test_ipc_memory_fallback' tests/integration/test_lists 2>/dev/null || true
printf '%s\n' '--- candidate test-list files ---'
fd -i '.*(yml|yaml|sqlite)$' tests/integration/test_lists 2>/dev/null | head -100Repository: NVIDIA/TensorRT-LLM
Length of output: 4859
Catch only AttributeError for a missing operator. PyTorch raises AttributeError when torch.ops.trtllm.allreduce is not registered. Broad exception handling can hide unrelated setup failures.
🧰 Tools
🪛 Ruff (0.16.1)
[warning] 50-50: Do not catch blind exception: Exception
(BLE001)
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
In `@tests/unittest/_torch/distributed/test_ipc_memory_fallback.py` around lines
46 - 51, Update _has_allreduce_op to catch only AttributeError when accessing
torch.ops.trtllm.allreduce, allowing unrelated setup failures to propagate while
preserving the false result for a missing operator.
Sources: Coding guidelines, Linters/SAST tools
| def cudaMalloc(self, size): | ||
| self.next_ptr += size | ||
| self.allocated.add(self.next_ptr) | ||
| return FakeCudaError.cudaSuccess, self.next_ptr | ||
|
|
||
| def cudaMemset(self, ptr, value, size): | ||
| return (FakeCudaError.cudaSuccess,) |
There was a problem hiding this comment.
🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win
🧩 Analysis chain
🏁 Script executed:
#!/bin/bash
set -euo pipefail
rg -n -C 3 \
'def cudaMalloc|def cudaMemset|test_.*(malloc|memset)|test_ipc_memory_fallback' \
tests/unittest/_torch/distributed/test_ipc_memory_fallback.py
rg -n -C 2 \
'test_ipc_memory_fallback|_torch/distributed' \
tests/integration/test_lists || trueRepository: NVIDIA/TensorRT-LLM
Length of output: 1175
🏁 Script executed:
#!/bin/bash
set -euo pipefail
printf '%s\n' '--- file outline ---'
ast-grep outline tests/unittest/_torch/distributed/test_ipc_memory_fallback.py
printf '%s\n' '--- test file ---'
cat -n tests/unittest/_torch/distributed/test_ipc_memory_fallback.py
printf '%s\n' '--- related source and test-list references ---'
rg -n -C 3 \
'cudaMalloc|cudaMemset|cudaSuccess|ipc.*fallback|NCCL|test_ipc_memory_fallback|unittest/_torch/distributed' \
tests/unittest/_torch/distributed tests/integration/test_lists
printf '%s\n' '--- applicable QA list files ---'
find tests/integration/test_lists -type f \( -path '*/qa/*' -o -iname '*qa*' \) -printRepository: NVIDIA/TensorRT-LLM
Length of output: 26052
🏁 Script executed:
#!/bin/bash
set -euo pipefail
printf '%s\n' '--- IpcMemory implementation locations ---'
rg -l 'class IpcMemory|def __init__.*IpcMemory|cudaMalloc|cudaMemset' --glob '*.py' .
printf '%s\n' '--- IpcMemory implementation ---'
rg -n -C 8 'class IpcMemory|cudaMalloc|cudaMemset|cudaIpcGetMemHandle|cudaIpcOpenMemHandle' \
--glob '*.py' .
printf '%s\n' '--- exact test registration across all integration lists ---'
rg -n -C 2 \
'test_ipc_memory_fallback\.py|unittest/_torch/distributed' \
tests/integration/test_lists
printf '%s\n' '--- test-list files and changed test paths ---'
git ls-files tests/integration/test_listsRepository: NVIDIA/TensorRT-LLM
Length of output: 38944
🏁 Script executed:
#!/bin/bash
set -euo pipefail
printf '%s\n' '--- implementation and error helper ---'
cat -n tensorrt_llm/_ipc_utils.py | sed -n '1,215p'
printf '%s\n' '--- applicable coding guidance ---'
find . -name CODING_GUIDELINES.md -print
if [ -f CODING_GUIDELINES.md ]; then
cat CODING_GUIDELINES.md
fi
printf '%s\n' '--- test-list guidance ---'
cat tests/integration/test_lists/test-db/README.md
printf '%s\n' '--- QA guidance ---'
cat tests/integration/test_lists/qa/README.md
printf '%s\n' '--- deterministic AST checks ---'
python3 - <<'PY'
import ast
from pathlib import Path
test_path = Path("tests/unittest/_torch/distributed/test_ipc_memory_fallback.py")
source_path = Path("tensorrt_llm/_ipc_utils.py")
test_tree = ast.parse(test_path.read_text())
test_names = [
node.name for node in test_tree.body
if isinstance(node, (ast.FunctionDef, ast.AsyncFunctionDef))
and node.name.startswith("test_")
]
print("test_functions:", test_names)
source_tree = ast.parse(source_path.read_text())
ipc_class = next(
node for node in source_tree.body
if isinstance(node, ast.ClassDef) and node.name == "IpcMemory"
)
init = next(node for node in ipc_class.body if isinstance(node, ast.FunctionDef)
and node.name == "__init__")
calls = []
for node in ast.walk(init):
if isinstance(node, ast.Call) and isinstance(node.func, ast.Attribute):
if node.func.attr in {"cudaMalloc", "cudaMemset", "_raise_if_error"}:
calls.append((node.lineno, node.func.attr))
print("IpcMemory.__init__ relevant calls:", sorted(calls))
registration = Path("tests/integration/test_lists/test-db/l0_cpu.yml").read_text()
print("directory_registration:",
"unittest/_torch/distributed" in registration)
qa_matches = []
for path in Path("tests/integration/test_lists/qa").rglob("*"):
if path.is_file() and path.name != ".gitignore":
text = path.read_text(errors="replace")
if "test_ipc_memory_fallback" in text:
qa_matches.append(str(path))
print("qa_matches:", qa_matches)
PYRepository: NVIDIA/TensorRT-LLM
Length of output: 50377
🏁 Script executed:
#!/bin/bash
set -euo pipefail
python3 - <<'PY'
import ast
from pathlib import Path
test_path = Path("tests/unittest/_torch/distributed/test_ipc_memory_fallback.py")
tree = ast.parse(test_path.read_text())
tests = [
node.name for node in tree.body
if isinstance(node, ast.FunctionDef) and node.name.startswith("test_")
]
print("test_functions:")
for name in tests:
print(f" {name}")
source = Path("tensorrt_llm/_ipc_utils.py").read_text()
for needle in (
"error, local_ptr = cudart.cudaMalloc(aligned_size)",
"_raise_if_error(error)",
"_raise_if_error(cudart.cudaMemset(local_ptr, 0, aligned_size)[0])",
):
print(f"{needle!r}: {needle in source}")
ci_list = Path("tests/integration/test_lists/test-db/l0_cpu.yml").read_text()
print(
"CI directory registration:",
"unittest/_torch/distributed" in ci_list,
)
qa_root = Path("tests/integration/test_lists/qa")
matches = [
str(path) for path in qa_root.rglob("*")
if path.is_file()
and path.name != ".gitignore"
and "test_ipc_memory_fallback" in path.read_text(errors="replace")
]
print("Manual-QA matches:", matches)
PYRepository: NVIDIA/TensorRT-LLM
Length of output: 655
Add CUDA allocation and memset failure tests.
Test coverage summary — insufficient. FakeCudart.cudaMalloc and FakeCudart.cudaMemset always succeed, so their RuntimeError paths are untested. Add configurable failures and tests that confirm errors propagate without IPC-to-NCCL fallback.
Added tests: test_ipc_ok, test_local_open_fails, test_peer_open_fails, test_local_export_fails, test_peer_export_fails, test_ipc_not_requested, test_inter_node_tp, test_mnnvl_kept_without_p2p, test_downgrade_to_nccl, test_moe_allreduce_raises, and test_lamport_skipped. The file is registered through unittest/_torch/distributed in tests/integration/test_lists/test-db/l0_cpu.yml; no matching manual-QA entry exists.
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
In `@tests/unittest/_torch/distributed/test_ipc_memory_fallback.py` around lines
92 - 98, Update FakeCudart.cudaMalloc and FakeCudart.cudaMemset to support
configurable CUDA failures, then add tests covering each RuntimeError path and
verifying errors propagate without falling back from IPC to NCCL. Keep the
existing success behavior unchanged and exercise the failure configuration
through the relevant distributed IPC test flows.
Source: Path instructions
|
@brnguyen2 all five points from your review are addressed and pushed, and the branch is rebased onto main (was 162 behind, now 0). the tests are CPU-safe now that #16498 moved them into |
|
/bot run |
|
PR_Github #68209 [ run ] triggered by Bot. Commit: |
|
PR_Github #68209 [ run ] completed with state
|
cudaDeviceCanAccessPeer reporting P2P support does not guarantee that CUDA IPC handles can be exported and imported. On 8x RTX 6000D, and on PCIe-only parts generally, cudaIpcOpenMemHandle fails with cudaErrorInvalidDevice and engine initialization aborted instead of using the NCCL path that already exists. IpcMemory.open_ipc_memory now returns None on such a failure, after agreeing on the outcome across the whole TP group and releasing whatever it managed to allocate, so that either every rank gets IPC buffers or none of them does. IpcMemory keeps null pointers in that case, and lamport_initialize is gated on the buffers actually being open. Fixes NVIDIA#16899 Signed-off-by: pjdurden <prajjwalchittori1@gmail.com>
Addresses review feedback on NVIDIA#17034. allreduce_workspace_has_ipc() was False for any workspace built without P2P, which includes ordinary inter-node TP, so it rewrote the strategy to NCCL before the MNNVL block and mnnvl_allreduce was never constructed on multi-node NVLink systems. IpcMemory now records ipc_failed separately from open_ipc, and only that failure mode, P2P reported but handles unusable, downgrades the strategy. Configurations that never had P2P keep the behaviour they had. get_allreduce_workspace returns the flag alongside the workspace, so a caller cannot read the thread-local state before it has been populated. The fallback is reported once, from open_ipc_memory where it is detected, and on every rank of the group rather than only on the ones that failed locally. AllReduce is constructed per decoder layer, so it no longer warns per instance. A rank that cannot export its handle contributes None to the allgather instead of an uninitialized cudaIpcMemHandle_t, so no rank imports a garbage handle. The number of collectives is unchanged. MoEAllReduce and MiniMaxAllReduceRMS have no NCCL path, so they now raise an actionable error instead of letting the kernel dereference null peer pointers. Tests cover the strategy downgrade, MNNVL still being selected when the workspace has no IPC buffers by design, and the lamport_initialize gating. Signed-off-by: pjdurden <prajjwalchittori1@gmail.com>
NVIDIA#16498 moved tests/unittest/_torch/distributed from l0_dgx_h100 to l0_cpu, so this file now runs with zero GPUs. The three tests that build an AllReduce skip when the trtllm custom ops are not registered, since AllReduce.__init__ resolves torch.ops.trtllm.allreduce before any of the logic under test. test_lamport_skipped still needs a GPU because the workspace helper it drives allocates device tensors. Signed-off-by: pjdurden <prajjwalchittori1@gmail.com>
|
@bo-nv thanks for running it. The pipeline links are all on the internal network so I cannot see which stage failed, could you share the failing stage or test name? In the meantime I rebased onto main, it was 484 behind, now 0, and the rebase was clean. Only one upstream commit had touched any of my files since my base, c0e6d79, and that one rewrites What I could check locally against the four changed files, all clean: ruff and ruff-format on Could you trigger another run on the rebased head, d0e8d04? If it fails again I would need the stage name to get any further. |
|
@bo-nv ping on this one. The pipeline links are internal so I still cannot see what broke. Could you paste the failing stage or test name? Branch is rebased and current, happy to turn a fix around quickly. |
Fall back to NCCL when CUDA IPC handles cannot be exchanged
Fixes #16899
1. Root cause
The reporter's follow-up (v1.3.0rc22) pins the failure precisely:
and notes that
can_access_peer()returnedTrueon the same machine.can_access_peer()(tensorrt_llm/_ipc_utils.py:39) only callscudaDeviceCanAccessPeer. That answers "can GPU A address GPU B's memory", whichis necessary but not sufficient for CUDA IPC: exporting and importing an IPC
memory handle is a separate capability that also fails on GPUs without CUDA IPC
support and when the ranks do not share an IPC namespace. On this system
cudaDeviceCanAccessPeersucceeds whilecudaIpcOpenMemHandledoes not.Because
can_access_peer()saidTrue,IpcMemorywas constructed withopen_ipc=True, andopen_ipc_memory()passed the raw CUDA error straight to_raise_if_error(). That unhandledRuntimeErroraborts engine initialization.This is a hard crash for a situation the runtime already fully supports. The
non-IPC path exists and works:
IpcMemory(open_ipc=False)leaves all pointersnull, and on the C++ side
AllReduceOp::ifFallbackToNCCL(
cpp/tensorrt_llm/thop/allreduceOp.cpp:1436) already routes every all-reduce toNCCL when P2P/NVLink is absent, which is exactly this hardware (RTX 6000D is a
PCIe part with no NVLink, so
mIsNVLINKSupportedis false and the custom kernelswould never have been selected anyway). Only the eager Python-side workspace
allocation stood in the way.
One secondary note that matches the report:
TRTLLM_DISABLE_CUSTOM_ALLREDUCE=1had no effect because no such env var exists anywhere in the tree (
grepreturnsnothing). The supported knob is
allreduce_strategy: NCCLin the LLM args, whichskips the workspace allocation entirely. That is a valid workaround, but users
should not need it.
2. The fix
Degrade to the already-supported NCCL path instead of crashing.
tensorrt_llm/_ipc_utils.py.IpcMemory.open_ipc_memoryno longer raiseswhen IPC handles cannot be exchanged. It returns
None, andIpcMemory.__init__reacts by setting
open_ipc = Falseandipc_failed = True, leaving thepointers null. Three details matter:
either every rank gets IPC buffers or none does. Without this, a rank that
succeeded would keep using buffers its peers never mapped, and the subsequent
tp_allgathercalls could desync.Noneto thehandle allgather rather than an uninitialized
cudaIpcMemHandle_t, so no rankever passes an uninitialized handle to
cudaIpcOpenMemHandle. That onecollective carries both the payload and the agreement on the export step, so
the total is still two collectives, as before.
is freed, so the fallback does not leak device memory.
cudaMalloc/cudaMemsetfailures still raise. Those are genuine errors, not acapability gap.
tensorrt_llm/_torch/distributed/allreduce_helper.py. Gatelamport_initialize()onlamport_buffers.open_ipcrather than onis_p2p_supported. Those two can now disagree, and callinglamport_initializewith a null
local_ptrwould be an illegal memory access.tensorrt_llm/_torch/distributed/ops.py.get_allreduce_workspacereturnsthe workspace together with an
ipc_failedflag, andAllReduce.__init__downgrades the strategy to
NCCLonly when that flag is set. The flagdistinguishes the new failure mode (P2P reported, IPC handles unusable) from a
workspace that holds no IPC buffers by design, which is the ordinary inter-node
TP case. That distinction is load-bearing: reacting to the latter would rewrite
the strategy before the MNNVL block and silently disable
mnnvl_allreduceonmulti-node NVLink systems. The
LOWPRECISIONworkspace allocation moved belowthe check so
initialize_static_lowprecision_buffersis not called on a nullworkspace.
SYMM_MEMis unaffected:AllReduce.forwardtriesself.symm_mem_allreducebefore consulting
self.strategy, so the downgrade only changes the fallbackpath.
Fused ops that have no NCCL path.
MoEAllReduceandMiniMaxAllReduceRMShand the workspace to custom kernels that reinterpret it as
void**, and neitherhas a non-IPC alternative, so the strategy downgrade above does not cover them.
They now raise an explicit error when
ipc_failedis set, rather thandereferencing the null peer pointers inside the kernel. That is the same
init-time failure those paths had before this PR, with a message that says why.
3. Files changed
tensorrt_llm/_ipc_utils.pyopen_ipc_memoryreturnsOptional[...]; collective IPC-failure agreement without exchanging uninitialized handles; cleanup;IpcMemory.ipc_failed; single warning on the grouptensorrt_llm/_torch/distributed/allreduce_helper.pylamport_initializeon the actual IPC statetensorrt_llm/_torch/distributed/ops.pyget_allreduce_workspacereturns(workspace, ipc_failed); strategy downgrade to NCCL on real IPC failure only; reorder LOWPRECISION allocation; explicit error forMoEAllReduce/MiniMaxAllReduceRMStests/unittest/_torch/distributed/test_ipc_memory_fallback.py4. Reported validation on a second platform
The reporter of #16899 applied the first revision of this patch on 8x RTX 5090
(PCIe, no NVLink), which is a different platform from the RTX 6000D in the
original report:
cudaErrorInvalidDevice: 101fromcudaIpcOpenMemHandle, on all 8 ranks.benchmark run to completion.
So the failure is not specific to the RTX 6000D, it affects PCIe-only parts
generally. That run predates the changes in section 2 that narrow the downgrade
and rework the handle exchange.
5. Risk and uncertainty
environment has neither GPUs nor
torch. The root cause is derived from thereporter's traceback plus reading the code; the specific reason
cudaIpcOpenMemHandlereturns 101 on that GPU is not established. The fix isdeliberately agnostic to that reason, it handles "IPC handles cannot be
exchanged" however it arises.
slower but correct NCCL path. That is the right trade-off here, since the
runtime already treats non-P2P topologies this way, but it does mean a
genuinely broken IPC setup degrades instead of failing loudly. The warning is
emitted once per process, from
open_ipc_memorywhere the condition isdetected, and on every rank of the group rather than only on the ones that
failed locally.
hit this initialization path; I have no evidence about the mechanism and make
no claim about it.
allgather that already existed, plus the one that covers the open step. Two
collectives per
IpcMemoryconstruction, unchanged.can_access_peer()still reportsTrueonsuch systems, so
self.is_p2p_supportedinmodeling_deepseekv3.py,modeling_glm.py,modeling_qwen3_moe.pyand friends remains optimistic. Acomplete fix would also plumb the IPC verdict into
AllReduceOp::setGroupTopology(cpp/tensorrt_llm/thop/allreduceOp.cpp) sothat C++ and Python agree from the start; the strategy downgrade above covers
the same ground from the Python side without touching C++.
concurrency 1) is a different failure from the rc22 init crash addressed here.
Their rc22 report supersedes it, and that is what this fix targets.
6. How I verified it
Constraints: no GPU and no
torchin this environment, sopytestcannot importtensorrt_llm. The pytest file was not executed. Instead:committed) loads the actual
tensorrt_llm/_ipc_utils.pyviaimportlibwiththe CUDA runtime,
Distributed,loggerandMappingstubbed, and runs sixscenarios. All assertions pass:
open_ipc=True,ipc_failed=False, non-null pointers, onehandle opened, two collectives, no warning.
cudaIpcOpenMemHandlereturnscudaErrorInvalidDevice(101), the reportedfailure: no exception,
ipc_failed=True, all pointers null, local bufferfreed, warning names the failing call.
opened handles closed, buffer freed, and it still warns rather than
degrading silently.
cudaIpcGetMemHandlefails locally: no handle is ever opened, buffer freed.cudaIpcOpenMemHandleis never called, so nouninitialized handle reaches the driver.
open_ipc=Falsewithipc_failed=False, no allocation, no collective. This is the case that mustnot disturb strategy selection.
ruffpinnedto the
v0.9.4this repo's pre-commit uses:ruff checkandruff format --diffclean on_ipc_utils.py,allreduce_helper.pyand the new test (non-legacy files).yapf --diffandisort --diffclean onops.py, which is inlegacy-files.txt, andscripts/legacy_utils.py lint-precommitreportsnothing on it.
Not verified: end-to-end
trtllm-bench --tp 8on affected hardware, and thatNCCL throughput is acceptable for this workload. Both need the reporter's
machine. The new tests are in
tests/unittest/_torch/distributed, which #16498moved from
l0_dgx_h100.ymlto the CPU-onlyl0_cpu.ymlon 2026-08-04, so CIruns them with zero GPUs. Everything there passes without a GPU; the three tests
that build an
AllReduceskip if the trtllm custom ops are not registered, andtest_lamport_skippedskips because the workspace helper it drives allocatesdevice tensors. The CUDA stubbing is
unittest.mock.patchscoped to a contextmanager and touches no real CUDA or MPI state, matching how
test_safe_mpi_comm.pyin the same directory already patches.Rebased onto
mainon 2026-08-05 (was 162 behind, now 0). Clean rebase: none ofthe three source files had upstream commits since the branch point.
Dev Engineer Review
AllReducefalls back to NCCL only after a genuine IPC exchange failure.get_allreduce_workspacereturn type is consistent with its callers.QA Engineer Review
tests/unittest/_torch/distributed/test_ipc_memory_fallback.py.tests/integration/test_lists/,test-db/, orqa/.