Skip to content

[https://nvbugs/6525010][fix] Fix TestLlama3_3_70BInstruct::test_nvfp… - #17040

Open
liji-nv wants to merge 5 commits into
NVIDIA:mainfrom
liji-nv:fix/nvbug6525010-userbuffer-release
Open

[https://nvbugs/6525010][fix] Fix TestLlama3_3_70BInstruct::test_nvfp…#17040
liji-nv wants to merge 5 commits into
NVIDIA:mainfrom
liji-nv:fix/nvbug6525010-userbuffer-release

Conversation

@liji-nv

@liji-nv liji-nv commented Jul 30, 2026

Copy link
Copy Markdown
Collaborator

…4_tp4 UB reuse

The B200 full accuracy sequence constructs several PyTorch engines in one process. UserBuffersManager updated its configured allocation size for each engine but retained CUDA VMM regions registered by earlier engines because UserBufferAllocator::deallocate was a no-op.

This allowed the Llama 3.1 8B TP4 case to leave a 64 MiB userbuffer in the singleton pool and the following Llama 3.3 70B TP4 case to reuse it for a 128 MiB request. The undersized mapping produced an invalid CUDA access that surfaced synchronously at the explicit graph-output copy when CUDA_LAUNCH_BLOCKING was enabled.

Add collective userbuffer unregistration that synchronizes ranks, unbinds multicast mappings, unmaps and frees virtual address ranges, releases local and imported allocation handles, and clears communicator metadata. Track logical releases independently so callers may release in any order while physical registrations are retired in the required LIFO order.

Expose UserBuffersManager shutdown through the Python binding and invoke it after model tensors and CUDA graphs are released during engine cleanup. Reject manager reinitialization while an earlier engine pool remains, preventing silent cross-engine reuse.

Extend the multi-GPU allocation tests to shut down each manager instance and repeatedly reinitialize it more than MAX_REGIONS times before changing allocation size. This verifies that shutdown unregisters physical regions instead of only dropping pool bookkeeping.

Remove the B200 waive for TestLlama3_3_70BInstruct::test_nvfp4_tp4[torch_compile=True]. The exact four-case sequence passed twice in one sbatch allocation with CUDA_LAUNCH_BLOCKING=1 (Slurm job 1554986), and the new userbuffer lifecycle regression passed in both iterations.

Dev Engineer Review

  • Added collective userbuffer unregistration with synchronization and complete CUDA VMM, multicast, allocation-handle, and communicator cleanup.
  • Added released-state tracking and stricter validation in UserBufferAllocator.
  • Added manager shutdown and reinitialization guards.
  • Moved collective userbuffer cleanup to deterministic executor shutdown.
  • Preserved managers during temporary shutdown.
  • Released CUDA graphs before userbuffer teardown.
  • Aggregated cleanup errors after all engines complete shutdown.
  • Exposed shutdown_userbuffers_manager() through nanobind.
  • Removed the B200 waiver for TestLlama3_3_70BInstruct::test_nvfp4_tp4[torch_compile=True].

QA Engineer Review

  • Added test_userbuffers_manager_reinitialize and _run_userbuffers_manager_reinitialize.
  • Added shutdown and cleanup tests in test_pytorch_model_engine.py.
  • Added executor shutdown tests in test_py_executor.py.
  • The new test functions are not listed in tests/integration/test_lists/test-db/ or tests/integration/test_lists/qa/.
  • Removed the B200 waiver entry from tests/integration/test_lists/waives.txt.

Verdict: needs follow-up (CBTS coverage mapping is unavailable for the new regression tests).

Description

Test Coverage

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-compatible or api-breaking. For api-breaking, include BREAKING in 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.

…4_tp4 UB reuse

The B200 full accuracy sequence constructs several PyTorch engines in one process. UserBuffersManager updated its configured allocation size for each engine but retained CUDA VMM regions registered by earlier engines because UserBufferAllocator::deallocate was a no-op.

This allowed the Llama 3.1 8B TP4 case to leave a 64 MiB userbuffer in the singleton pool and the following Llama 3.3 70B TP4 case to reuse it for a 128 MiB request. The undersized mapping produced an invalid CUDA access that surfaced synchronously at the explicit graph-output copy when CUDA_LAUNCH_BLOCKING was enabled.

Add collective userbuffer unregistration that synchronizes ranks, unbinds multicast mappings, unmaps and frees virtual address ranges, releases local and imported allocation handles, and clears communicator metadata. Track logical releases independently so callers may release in any order while physical registrations are retired in the required LIFO order.

Expose UserBuffersManager shutdown through the Python binding and invoke it after model tensors and CUDA graphs are released during engine cleanup. Reject manager reinitialization while an earlier engine pool remains, preventing silent cross-engine reuse.

Extend the multi-GPU allocation tests to shut down each manager instance and repeatedly reinitialize it more than MAX_REGIONS times before changing allocation size. This verifies that shutdown unregisters physical regions instead of only dropping pool bookkeeping.

Remove the B200 waive for TestLlama3_3_70BInstruct::test_nvfp4_tp4[torch_compile=True]. The exact four-case sequence passed twice in one sbatch allocation with CUDA_LAUNCH_BLOCKING=1 (Slurm job 1554986), and the new userbuffer lifecycle regression passed in both iterations.

Signed-off-by: Jin Li <59594262+liji-nv@users.noreply.github.com>
@liji-nv

liji-nv commented Jul 30, 2026

Copy link
Copy Markdown
Collaborator Author

/bot run --disable-fail-fast

@tensorrt-cicd

Copy link
Copy Markdown
Collaborator

PR_Github #62665 [ run ] triggered by Bot. Commit: 58e8399 Link to invocation

@coderabbitai

coderabbitai Bot commented Jul 30, 2026

Copy link
Copy Markdown
Contributor

Review Change Stack

Note

Reviews paused

It 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 reviews.auto_review.auto_pause_after_reviewed_commits setting.

Use the following commands to manage reviews:

  • @coderabbitai resume to resume automatic reviews.
  • @coderabbitai review to trigger a single review.

Use the checkboxes below for quick actions:

  • ▶️ Resume reviews
  • 🔍 Trigger review

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 5a4c24fa-ae18-4f60-a1a2-98e02f77f23d

📥 Commits

Reviewing files that changed from the base of the PR and between 3ec5cbf and b2b1597.

📒 Files selected for processing (5)
  • tensorrt_llm/_torch/pyexecutor/_util.py
  • tensorrt_llm/_torch/pyexecutor/model_engine.py
  • tensorrt_llm/_torch/pyexecutor/py_executor.py
  • tests/integration/test_lists/waives.txt
  • tests/unittest/_torch/executor/test_pytorch_model_engine.py
💤 Files with no reviewable changes (1)
  • tests/integration/test_lists/waives.txt
🚧 Files skipped from review as they are similar to previous changes (4)
  • tensorrt_llm/_torch/pyexecutor/_util.py
  • tests/unittest/_torch/executor/test_pytorch_model_engine.py
  • tensorrt_llm/_torch/pyexecutor/py_executor.py
  • tensorrt_llm/_torch/pyexecutor/model_engine.py

Walkthrough

User-buffer allocation and collective registration now support explicit release. A manager shutdown API is exposed to Python and integrated into model-engine and executor shutdown. Tests cover cleanup, repeated reinitialization, region reuse, and shutdown failures.

Changes

User-buffer lifecycle

Layer / File(s) Summary
Allocator release tracking
cpp/tensorrt_llm/kernels/userbuffers/ub_allocator.*
Allocator state tracks released buffers, implements deallocation, removes released entries, and rejects invalid or released accesses.
Collective buffer unregistration
cpp/tensorrt_llm/kernels/userbuffers/userbuffers.*
Registration bounds checks and collective unregistration release CUDA, multicast, communicator, and host-side resources.
Manager shutdown API
cpp/tensorrt_llm/kernels/userbuffers/userbuffersManager.*, cpp/tensorrt_llm/nanobind/userbuffers/bindings.cpp
The manager adds shutdown and initialization guards. Python bindings expose manager shutdown.
Engine shutdown integration
tensorrt_llm/_torch/pyexecutor/model_engine.py, tensorrt_llm/_torch/pyexecutor/py_executor.py, tensorrt_llm/_torch/pyexecutor/encoder_executor.py, tensorrt_llm/_torch/pyexecutor/_util.py
Engine and executor shutdown paths control collective user-buffer teardown, preserve managers when requested, and report teardown failures.
Lifecycle validation
tests/unittest/_torch/multi_gpu/test_allocate_output_buffer_kinds.py, tests/unittest/_torch/executor/test_pytorch_model_engine.py, tests/unittest/_torch/executor/test_py_executor.py, tests/integration/test_lists/waives.txt
Tests cover shutdown, failure handling, repeated reinitialization, region reuse, and waiver removal.

Estimated code review effort: 4 (Complex) | ~45 minutes

Suggested labels: api-compatible

Suggested reviewers: brnguyen2, pcastongue

Sequence Diagram(s)

sequenceDiagram
  participant PyExecutor
  participant PyTorchModelEngine
  participant UserBuffersManager
  participant UserBufferAllocator
  PyExecutor->>PyTorchModelEngine: request user-buffer shutdown
  PyTorchModelEngine->>UserBuffersManager: shut down initialized manager
  UserBuffersManager->>UserBufferAllocator: deallocate tracked buffers
  UserBufferAllocator-->>UserBuffersManager: release allocator entries
  UserBuffersManager-->>PyTorchModelEngine: return teardown status
  PyTorchModelEngine-->>PyExecutor: report shutdown result
Loading
🚥 Pre-merge checks | ✅ 4 | ❌ 1

❌ Failed checks (1 warning)

Check name Status Explanation Resolution
Docstring Coverage ⚠️ Warning Docstring coverage is 39.13% which is insufficient. The required threshold is 80.00%. Write docstrings for the functions missing them to satisfy the coverage threshold.
✅ Passed checks (4 passed)
Check name Status Explanation
Title check ✅ Passed The title follows the required NVBugs and fix format and identifies the failing test addressed by the userbuffer lifecycle changes.
Description check ✅ Passed The description clearly explains the issue, solution, lifecycle changes, test coverage, validation results, and checklist completion.
Linked Issues check ✅ Passed Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check ✅ Passed Check skipped because no linked issues were found for this pull request.
✨ Finishing Touches
🧪 Generate unit tests (beta)
  • Create PR with unit tests

Comment @coderabbitai help to get the list of available commands.

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actionable comments posted: 1

🧹 Nitpick comments (2)
cpp/tensorrt_llm/kernels/userbuffers/userbuffers-host.cpp (1)

436-483: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

Teardown correctly mirrors registration in reverse; strict LIFO invariant enforced via TLLM_CHECK.

The unmap/free/release sequence, the mc_offset rewind (with an assertion tying it back to the recorded mc_ptr[handle]), the host free() of uchandles/peer_ptr, and the cudaMemset clearing the peer-pointer table all correctly reverse the corresponding steps in register_user_buffer_collective. This function is collective (an initial cudaDeviceSynchronize/ub_barrier, and a trailing ub_barrier), so all ranks must call it in lockstep — see the related comment on UserBuffersManager::shutdown() for a cross-rank ordering concern this introduces.

Consider adding a short comment above this function documenting that callers must release buffers strictly in reverse-allocation order (enforced by TLLM_CHECK(handle == comm->free_region - 1)), since that invariant is easy to miss for future callers of this low-level API.

🤖 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 `@cpp/tensorrt_llm/kernels/userbuffers/userbuffers-host.cpp` around lines 436 -
483, Add a brief comment immediately above unregister_user_buffer_collective
documenting that buffers must be released in strict reverse-allocation order.
Reference the existing handle == comm->free_region - 1 TLLM_CHECK as the
enforcement of this LIFO invariant, without changing teardown behavior.
tensorrt_llm/_torch/pyexecutor/model_engine.py (1)

2625-2698: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

Cleanup ordering and idempotent-retry design look correct.

Releasing cuda graphs / model refs / legacy ub_buffers before release_gc() + conditional shutdown_userbuffers_manager() is the right order, and leaving _userbuffers_manager_initialized / _cleanup_done unset on failure correctly enables a safe retry on the next cleanup() call.

Minor: the docstring's "Raises" section only calls out the RuntimeError chained from failed legacy userbuffer deallocations; ub.shutdown_userbuffers_manager() (line 2696) can independently raise (e.g. if a manager-tracked buffer is still marked in-use) and isn't mentioned. Worth a one-line addition for future maintainers debugging a raised exception here.

🤖 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/pyexecutor/model_engine.py` around lines 2625 - 2698,
Update the cleanup() docstring’s Raises section to also document that
ub.shutdown_userbuffers_manager() may raise independently, including when
manager-tracked buffers remain in use; keep the existing legacy userbuffer
deallocation RuntimeError description unchanged.
🤖 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 `@cpp/tensorrt_llm/kernels/userbuffers/userbuffersManager.cpp`:
- Around line 84-98: Ensure PyTorchModelEngine.cleanup() invokes
ub.shutdown_userbuffers_manager() on every rank even when an earlier legacy
ub_buffers cleanup operation raises. Preserve the collective teardown ordering
by capturing or deferring the exception, completing manager shutdown across all
ranks, then propagating the cleanup error afterward.

---

Nitpick comments:
In `@cpp/tensorrt_llm/kernels/userbuffers/userbuffers-host.cpp`:
- Around line 436-483: Add a brief comment immediately above
unregister_user_buffer_collective documenting that buffers must be released in
strict reverse-allocation order. Reference the existing handle ==
comm->free_region - 1 TLLM_CHECK as the enforcement of this LIFO invariant,
without changing teardown behavior.

In `@tensorrt_llm/_torch/pyexecutor/model_engine.py`:
- Around line 2625-2698: Update the cleanup() docstring’s Raises section to also
document that ub.shutdown_userbuffers_manager() may raise independently,
including when manager-tracked buffers remain in use; keep the existing legacy
userbuffer deallocation RuntimeError description unchanged.
🪄 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: e147c6f8-859e-4922-9d23-7b0f5ad908be

📥 Commits

Reviewing files that changed from the base of the PR and between 8624ec9 and 58e8399.

📒 Files selected for processing (10)
  • cpp/tensorrt_llm/kernels/userbuffers/ub_allocator.cpp
  • cpp/tensorrt_llm/kernels/userbuffers/ub_allocator.h
  • cpp/tensorrt_llm/kernels/userbuffers/userbuffers-host.cpp
  • cpp/tensorrt_llm/kernels/userbuffers/userbuffers.h
  • cpp/tensorrt_llm/kernels/userbuffers/userbuffersManager.cpp
  • cpp/tensorrt_llm/kernels/userbuffers/userbuffersManager.h
  • cpp/tensorrt_llm/nanobind/userbuffers/bindings.cpp
  • tensorrt_llm/_torch/pyexecutor/model_engine.py
  • tests/integration/test_lists/waives.txt
  • tests/unittest/_torch/multi_gpu/test_allocate_output_buffer_kinds.py
💤 Files with no reviewable changes (1)
  • tests/integration/test_lists/waives.txt

Comment thread cpp/tensorrt_llm/kernels/userbuffers/userbuffersManager.cpp
… shutdown

PyTorchModelEngine.cleanup raised immediately when a legacy directly-owned userbuffer failed to deallocate. On a multi-rank engine, that rank skipped UserBuffersManager shutdown while peers could enter its collective unregister sequence, risking mismatched teardown or a hang.

Collect legacy deallocation failures, release tensor-owned resources, and invoke shutdown_userbuffers_manager on every initialized rank before propagating the original cleanup error. Preserve failed legacy buffers and the incomplete cleanup state so deterministic callers can retry safely.

Add a focused cleanup regression that injects a legacy deallocation failure and verifies manager shutdown still runs before the error is returned, while the failed buffer and retry state are retained.

Signed-off-by: Jin Li <59594262+liji-nv@users.noreply.github.com>
@liji-nv

liji-nv commented Jul 30, 2026

Copy link
Copy Markdown
Collaborator Author

/bot run --disable-fail-fast

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🧹 Nitpick comments (1)
tests/unittest/_torch/executor/test_pytorch_model_engine.py (1)

153-181: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

Exercise the retry completion path.

Coverage summary: added PyTorchModelEngineTestCase.test_cleanup_shuts_down_manager_after_legacy_userbuffer_error; no tests were modified or removed. The test covers deferred errors, shutdown, flag reset, and failed-buffer retention, but the coverage verdict is insufficient because it never calls cleanup() again to verify the failed buffer is retried and _cleanup_done becomes True.

As per path instructions, test-code changes must include a coverage summary and verdict.

Suggested extension
-        with patch(
+        with patch(
                 "tensorrt_llm._torch.pyexecutor.model_engine.ub.ub_deallocate",
-                side_effect=legacy_error), patch(
+                side_effect=[legacy_error, None]) as deallocate, patch(
                     "tensorrt_llm._torch.pyexecutor.model_engine."
                     "ub.shutdown_userbuffers_manager") as shutdown_manager, \
                 patch("tensorrt_llm._torch.pyexecutor.model_engine.release_gc"):
             with self.assertRaisesRegex(
                     RuntimeError,
                     "Failed to deallocate one or more userbuffers"):
                 engine.cleanup()
+            engine.cleanup()

+        self.assertEqual(deallocate.call_count, 2)
         shutdown_manager.assert_called_once_with()
+        self.assertIsNone(engine.ub_buffers)
         self.assertFalse(engine._userbuffers_manager_initialized)
-        self.assertEqual(len(engine.ub_buffers), 1)
-        self.assertFalse(engine._cleanup_done)
+        self.assertTrue(engine._cleanup_done)
🤖 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/executor/test_pytorch_model_engine.py` around lines 153
- 181, Extend test_cleanup_shuts_down_manager_after_legacy_userbuffer_error to
call engine.cleanup() a second time after the expected first-call exception.
Verify the retained userbuffer is retried successfully, _cleanup_done becomes
True, and the manager shutdown behavior remains correct; update the test’s
coverage summary and verdict to reflect the completed retry-path coverage.

Source: 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.

Nitpick comments:
In `@tests/unittest/_torch/executor/test_pytorch_model_engine.py`:
- Around line 153-181: Extend
test_cleanup_shuts_down_manager_after_legacy_userbuffer_error to call
engine.cleanup() a second time after the expected first-call exception. Verify
the retained userbuffer is retried successfully, _cleanup_done becomes True, and
the manager shutdown behavior remains correct; update the test’s coverage
summary and verdict to reflect the completed retry-path coverage.

ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: e482eb37-b5e9-4507-85ab-f84f54a943b3

📥 Commits

Reviewing files that changed from the base of the PR and between 58e8399 and 4da3d48.

📒 Files selected for processing (2)
  • tensorrt_llm/_torch/pyexecutor/model_engine.py
  • tests/unittest/_torch/executor/test_pytorch_model_engine.py

@tensorrt-cicd

Copy link
Copy Markdown
Collaborator

PR_Github #62729 [ run ] triggered by Bot. Commit: 4da3d48 Link to invocation

@tensorrt-cicd

Copy link
Copy Markdown
Collaborator

PR_Github #62665 [ run ] completed with state ABORTED. Commit: 58e8399

Link to invocation

@BowenFu

BowenFu commented Jul 30, 2026

Copy link
Copy Markdown

The unregister mechanics look careful — the LIFO drain via mReleased handles out-of-order release cleanly, and the free_region > MAX_REGIONS>= change is a real out-of-bounds fix (the arrays are [MAX_REGIONS] with MAX_REGIONS 16, so the old > let free_region == 16 write index 16).

Two concerns, both about callers, and both stemming from ub_deallocate no longer being a no-op.

1. ub_deallocate becomes collective, and it's reachable from __del__.

unregister_user_buffer_collective does cudaDeviceSynchronize() + ub_barrier(comm->comm_intra) before the unmap and again after, and ub_barrier is MPI_Barrier (userbuffers-host.cpp:115-118). Today UserBufferAllocator::deallocate is {}, so ub.ub_deallocate is a pure no-op on every path.

ModelEngine.cleanup() calls ub.ub_deallocate(u.addr) per buffer (model_engine.py:2335), and its own docstring says it runs from PyExecutor.shutdown() and from __del__ (:2352-2366), "best-effort fallback during garbage collection / interpreter shutdown". GC timing is per-rank and non-deterministic, and __del__ may not run at all at interpreter shutdown. After this change, a teardown reaching cleanup() through __del__ can have ranks entering MPI_Barrier a different number of times, in a different order, or not at all — that's a hang, or an MPI call after MPI_Finalize.

Worth either confirming that path is collective-safe, or restricting the actual unregister to the deterministic shutdown() path.

2. The new LIFO precondition contradicts the retry contract cleanup() documents.

deallocate() sets mReleased[index] = true and then drains. If a drain step throws, the mark is already set. cleanup() deliberately keeps failed buffers attached "so a deterministic cleanup() call can retry without double-freeing buffers that were already released" — but that retry calls deallocate(addr) again and hits TLLM_CHECK(!mReleased[index]), which aborts instead of retrying.

Non-blocking: TLLM_CHECK(handle > 0) reads like an off-by-one until you find userbuffers-host.cpp:286 reserving handle 0 for comm->gpu_ptrs. It's correct as written — maybe worth a one-line comment saying so.

@tensorrt-cicd

Copy link
Copy Markdown
Collaborator

PR_Github #62729 [ run ] completed with state SUCCESS. Commit: 4da3d48
/LLM/main/L0_MergeRequest_PR pipeline #50863 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

@fredricz-20070104 fredricz-20070104 left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Review summary - CONCERNS

Verdict: The core fix is correct and directly addresses the no-op deallocate root cause of nvbugs/6525010, but the new teardown runs cross-rank collective barriers whose failure modes are not covered by tests, so it merits confirmation before merge.

Concerns

  1. [MAJOR] cpp/tensorrt_llm/kernels/userbuffers/userbuffers-host.cpp:452 (and userbuffersManager.cpp shutdown()) - collective teardown can hang or crash across ranks
    • What is wrong: unregister_user_buffer_collective does cudaDeviceSynchronize + ub_barrier(comm->comm_intra) at entry and a trailing ub_barrier. UserBuffersManager::shutdown() first asserts !buffer.second ("still in use") for every buffer via TLLM_CHECK_WITH_INFO before any teardown.
    • How it fails: these are collective ops requiring every rank in lockstep and identical LIFO order. If one rank still has a buffer in-use (or releases in a different order) while a peer does not, the peers pass the check and block on ub_barrier waiting for the rank that threw before the barrier — a hang. Any TLLM_CHECK failure inside unregister on one rank aborts it mid-collective while peers wait forever.
    • Suggested fix: guarantee cleanup()/shutdown() is invoked symmetrically on all ranks with identical buffer state, and validate the in-use invariant before any rank enters a barrier (ideally collectively), so no rank can abort after a peer has entered the barrier.

Minor notes (non-blocking)

  • cpp/tensorrt_llm/kernels/userbuffers/userbuffers-host.cpp:439 - TLLM_CHECK(handle > 0) forbids unregistering handle 0; if a real userbuffer ever lands on handle 0 the LIFO pop aborts. Document that region 0 is reserved, or use handle >= 0.
  • tests/unittest/_torch/executor/test_pytorch_model_engine.py:153 - test never calls cleanup() a second time, so the retried-buffer completion path (ub_buffers cleared, _cleanup_done True) is untested.
  • tests/unittest/_torch/multi_gpu/test_allocate_output_buffer_kinds.py:205 - reinit test returns True when ub_supported() is False; use pytest.skip so a non-run is not reported as a pass.

QA view

  • Test coverage: partial - the reinit multi-GPU test exercises shutdown/reinit past MAX_REGIONS, and the unit test covers deferred-error ordering, but the retry-completion path and the shutdown-while-in-use abort path are uncovered.
  • SM coverage: the userbuffer VMM/multicast teardown is architecture-independent, so the 2-GPU unit test is a fair functional check; the arch-specific nvfp4/B200 failure is validated only by the un-waived TestLlama3_3_70BInstruct::test_nvfp4_tp4[torch_compile=True] running on B200 in CI, which cannot be confirmed from the diff.
  • Test code: reinit test silently passes on unsupported HW; single-tensor-per-iteration does not exercise multi-buffer LIFO within one manager instance.
  • Test time: significant - removing the waive re-enables a full B200 70B nvfp4 accuracy case; the reinit test adds smaller multi-GPU time.
  • Needs /qa-verify: yes - a B200 waive is removed for an arch-specific fix; QA should confirm the un-waived accuracy test passes reliably on B200 in CI and that the shutdown/cleanup collective path does not hang.

Does this actually fix 6525010?

Yes. The bug is stale reuse of an undersized 64MiB VMM region because deallocate was a no-op; the fix makes deallocate call unregister_user_buffer_collective (LIFO), adds shutdown() to retire the pool, calls it during engine cleanup, and rejects re-init while a prior pool exists. The >>= MAX_REGIONS change additionally closes an off-by-one. I could not independently rerun the B200 sequence, so this rests on the un-waived integration test.

Possible new issues

  • Cross-rank deadlock/crash on asymmetric shutdown as described above.
  • mc_offset rewind assumes every torn-down region had UB_MEM_MC_CREATED; mixed MC/non-MC regions could desync the multicast offset.
  • The >= capacity change reduces registrable regions by one versus prior behaviour.

What I could not verify

  • Whether region 0 is reserved at init (relevant to the handle > 0 assertion).
  • Whether cleanup()/shutdown() is always called symmetrically across ranks with identical buffer state (the collective safety hinges on this).
  • Runtime behaviour of the un-waived B200 accuracy test in CI.

Automated review by NVCortex Lite, run by @fredricz-20070104.

…eardown

Userbuffer deallocation now performs CUDA synchronization and MPI barriers, so invoking it from rank-asynchronous destructor cleanup can hang or run after MPI finalization. It is also unsafe to retry after a buffer has entered the unregister drain because CUDA VMM teardown is not transactional.

Move userbuffer release into deterministic terminal executor shutdown. Release CUDA graphs first, synchronize the device, visit target and draft engines in a fixed order on every rank, and defer errors until all engine collectives have been attempted. Keep destructor cleanup non-collective and make failed unregister fail-stop instead of retrying partially released buffers.

Preserve userbuffer managers during the temporary executor shutdown used for KV-cache capacity estimation. The final executor reuses those engines, and prematurely shutting down the manager reset its capacity to zero, causing later TP4 allocations to fail. Encoder shutdown now follows the same explicit graph-before-userbuffer ordering.

Add focused coverage for deferred legacy cleanup errors, non-collective destructor cleanup, all-engine collective ordering, and nonterminal manager preservation. Improve the manager capacity assertion to report requested and available bytes.

Validated on one 4xB200 allocation with CUDA_LAUNCH_BLOCKING=1. Two consecutive iterations passed the focused unit tests, multi-GPU manager reinitialization test, and all four integration cases, including TestLlama3_3_70BInstruct::test_nvfp4_tp4[torch_compile=True]. Slurm job 1563971 completed with exit code 0.

Signed-off-by: Jin Li <59594262+liji-nv@users.noreply.github.com>

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actionable comments posted: 4

Caution

Some comments are outside the diff and can’t be posted inline due to platform limitations.

⚠️ Outside diff range comments (1)
tests/unittest/_torch/executor/test_pytorch_model_engine.py (1)

161-175: 🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win

Assert the cleanup order.

This test proves that shutdown_userbuffers_manager() runs before shutdown_userbuffers() returns. It does not prove that ub_deallocate() runs before manager shutdown.

Record both side effects and assert the deallocation call occurs before manager shutdown. This protects the deferred-error contract.

🤖 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/executor/test_pytorch_model_engine.py` around lines 161
- 175, Update the test around engine.shutdown_userbuffers() to record calls to
ub_deallocate and shutdown_userbuffers_manager in a shared order list, then
assert the deallocation call occurs before manager shutdown. Keep the existing
exception and state assertions, and ensure both mocked side effects append
distinct markers when invoked.
🤖 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/pyexecutor/encoder_executor.py`:
- Around line 59-68: Make EncoderExecutor.shutdown idempotent by adding a
completion guard, such as _cleanup_done, before accessing self.model_engine.
Return immediately on subsequent calls, and mark cleanup complete only after the
existing release, synchronization, and shutdown_userbuffers flow finishes while
preserving the finally-based model_engine deletion.

In `@tensorrt_llm/_torch/pyexecutor/model_engine.py`:
- Around line 2670-2679: Update the shutdown error handling around the
userbuffer deallocation and manager shutdown checks so that when both ub_errors
and manager_error are present, the raised RuntimeError includes diagnostic
details from both failures. Preserve the existing single-failure messages and
exception chaining behavior for cases where only one failure occurs, while
keeping _userbuffers_shutdown_failed set for any failure.

In `@tensorrt_llm/_torch/pyexecutor/py_executor.py`:
- Around line 1533-1537: Update the userbuffer shutdown error handling around
userbuffer_errors to log every collected failure, including errors beyond
userbuffer_errors[0], before raising the RuntimeError. Preserve the existing
aggregated exception and first-error chaining while ensuring each engine
shutdown failure is surfaced through the established logging mechanism.

In `@tests/unittest/_torch/executor/test_py_executor.py`:
- Around line 111-124: Annotate both new test functions with monkeypatch:
pytest.MonkeyPatch and -> None, and annotate the nested fail_target_userbuffers
function with -> None. Keep the existing test behavior and coverage unchanged.

---

Outside diff comments:
In `@tests/unittest/_torch/executor/test_pytorch_model_engine.py`:
- Around line 161-175: Update the test around engine.shutdown_userbuffers() to
record calls to ub_deallocate and shutdown_userbuffers_manager in a shared order
list, then assert the deallocation call occurs before manager shutdown. Keep the
existing exception and state assertions, and ensure both mocked side effects
append distinct markers when invoked.
🪄 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: 02e081bd-ddb3-4ed5-bfbd-fe8feb192c26

📥 Commits

Reviewing files that changed from the base of the PR and between 4da3d48 and 427b582.

📒 Files selected for processing (7)
  • cpp/tensorrt_llm/kernels/userbuffers/userbuffersManager.cpp
  • tensorrt_llm/_torch/pyexecutor/_util.py
  • tensorrt_llm/_torch/pyexecutor/encoder_executor.py
  • tensorrt_llm/_torch/pyexecutor/model_engine.py
  • tensorrt_llm/_torch/pyexecutor/py_executor.py
  • tests/unittest/_torch/executor/test_py_executor.py
  • tests/unittest/_torch/executor/test_pytorch_model_engine.py
🚧 Files skipped from review as they are similar to previous changes (1)
  • cpp/tensorrt_llm/kernels/userbuffers/userbuffersManager.cpp

Comment thread tensorrt_llm/_torch/pyexecutor/encoder_executor.py
Comment thread tensorrt_llm/_torch/pyexecutor/model_engine.py
Comment thread tensorrt_llm/_torch/pyexecutor/py_executor.py
Comment thread tests/unittest/_torch/executor/test_py_executor.py Outdated
@liji-nv

liji-nv commented Jul 31, 2026

Copy link
Copy Markdown
Collaborator Author

The unregister mechanics look careful — the LIFO drain via mReleased handles out-of-order release cleanly, and the free_region > MAX_REGIONS>= change is a real out-of-bounds fix (the arrays are [MAX_REGIONS] with MAX_REGIONS 16, so the old > let free_region == 16 write index 16).

Two concerns, both about callers, and both stemming from ub_deallocate no longer being a no-op.

1. ub_deallocate becomes collective, and it's reachable from __del__.

unregister_user_buffer_collective does cudaDeviceSynchronize() + ub_barrier(comm->comm_intra) before the unmap and again after, and ub_barrier is MPI_Barrier (userbuffers-host.cpp:115-118). Today UserBufferAllocator::deallocate is {}, so ub.ub_deallocate is a pure no-op on every path.

ModelEngine.cleanup() calls ub.ub_deallocate(u.addr) per buffer (model_engine.py:2335), and its own docstring says it runs from PyExecutor.shutdown() and from __del__ (:2352-2366), "best-effort fallback during garbage collection / interpreter shutdown". GC timing is per-rank and non-deterministic, and __del__ may not run at all at interpreter shutdown. After this change, a teardown reaching cleanup() through __del__ can have ranks entering MPI_Barrier a different number of times, in a different order, or not at all — that's a hang, or an MPI call after MPI_Finalize.

Worth either confirming that path is collective-safe, or restricting the actual unregister to the deterministic shutdown() path.

2. The new LIFO precondition contradicts the retry contract cleanup() documents.

deallocate() sets mReleased[index] = true and then drains. If a drain step throws, the mark is already set. cleanup() deliberately keeps failed buffers attached "so a deterministic cleanup() call can retry without double-freeing buffers that were already released" — but that retry calls deallocate(addr) again and hits TLLM_CHECK(!mReleased[index]), which aborts instead of retrying.

Non-blocking: TLLM_CHECK(handle > 0) reads like an off-by-one until you find userbuffers-host.cpp:286 reserving handle 0 for comm->gpu_ptrs. It's correct as written — maybe worth a one-line comment saying so.

1 has been fixed, the cleanup now runs explicitly in shutdown.

2 is not a real concern. We do not actually expect any retry. Any retry means the whole Buffer is not in good state.

@liji-nv

liji-nv commented Jul 31, 2026

Copy link
Copy Markdown
Collaborator Author

Concerns

  1. [MAJOR] cpp/tensorrt_llm/kernels/userbuffers/userbuffers-host.cpp:452 (and userbuffersManager.cpp shutdown()) - collective teardown can hang or crash across ranks

    • What is wrong: unregister_user_buffer_collective does cudaDeviceSynchronize + ub_barrier(comm->comm_intra) at entry and a trailing ub_barrier. UserBuffersManager::shutdown() first asserts !buffer.second ("still in use") for every buffer via TLLM_CHECK_WITH_INFO before any teardown.
    • How it fails: these are collective ops requiring every rank in lockstep and identical LIFO order. If one rank still has a buffer in-use (or releases in a different order) while a peer does not, the peers pass the check and block on ub_barrier waiting for the rank that threw before the barrier — a hang. Any TLLM_CHECK failure inside unregister on one rank aborts it mid-collective while peers wait forever.
    • Suggested fix: guarantee cleanup()/shutdown() is invoked symmetrically on all ranks with identical buffer state, and validate the in-use invariant before any rank enters a barrier (ideally collectively), so no rank can abort after a peer has entered the barrier.

We verified the current teardown path and believe rank state remains symmetric by construction:

  • UB region registration is collective, so allocation count, handles, and order must already match across ranks.
  • TP ranks execute the same compiled UB graph and use the same pool allocation/release order.
  • Terminal shutdown now runs deterministically on every rank, releases graphs first, synchronizes CUDA, and shuts down target/draft engines in a fixed order.
  • A failed unregister is fail-stop and cannot be retried, preventing subsequent state divergence.

Therefore, the manager’s in-use check should either fail on every rank before unregister begins or pass on every rank, after which all ranks drain the same LIFO sequence.

The described hang remains theoretically possible after an existing rank-local fatal condition, such as a CUDA error, process loss, asymmetric internal-API use, or another bug that has already violated the collective contract. A preflight collective would not handle a rank that never enters teardown and would add another collective without addressing general MPI fault tolerance.

Given that we found no supported execution path that can produce divergent manager state, we do not consider this an actionable blocker for the current change. Collective state validation could still be considered separately as defensive hardening.

…down

Make EncoderExecutor shutdown idempotent so repeated terminal cleanup does not access an already-deleted model engine. The completion guard is set only after graph release, device synchronization, and the userbuffer shutdown attempt have run, while retaining unconditional engine deletion.

Preserve complete diagnostics when teardown fails. Combined legacy-buffer and manager failures now report both details while chaining the first deallocation error, and PyExecutor logs every target or draft engine shutdown failure before raising its aggregate error.

Extend focused coverage for shutdown ordering, nonterminal manager preservation, repeated encoder shutdown, multi-engine error logging, and combined manager/deallocation failures. Add the requested pytest type annotations without changing existing test behavior.

Validation: all pre-commit hooks passed. On B200, Slurm job 1567861 completed with exit code 0 and all seven focused teardown tests passed.
Signed-off-by: Jin Li <59594262+liji-nv@users.noreply.github.com>
@BowenFu

BowenFu commented Jul 31, 2026

Copy link
Copy Markdown

Re-reviewed at 427b5823c. Both of my earlier findings are addressed, and in a better shape than I asked for:

  1. Collective teardown out of __del__shutdown_userbuffers() is now separate from cleanup(), documented as never-from-__del__, and PyExecutor.shutdown() enters it on every rank in a fixed engine order (main, then draft), collecting errors and raising only after the loop. That is the rank-symmetric ordering the MPI_Barrier in unregister_user_buffer_collective needs. shutdown_userbuffers=False for capacity estimation is the right escape hatch.
  2. Retry contract — instead of re-entering deallocate(addr) into TLLM_CHECK(!mReleased[index]), _userbuffers_shutdown_failed now makes a second call raise up front. Given VMM teardown isn't transactional, refusing the retry is more honest than pretending it's recoverable.

One residual thing, non-blocking: shutdown_userbuffers() early-returns when not ub_buffers and not manager_initialized, and that test is rank-local. If any rank in the TP group ever reaches shutdown with a different buffer/manager state than its peers, the peers enter unregister_user_buffer_collective's barrier and that rank doesn't — a hang rather than an error. I think it's unreachable today because _init_userbuffers is config-driven and uniform across TP ranks, but the invariant "every rank has the same number of UB buffers at shutdown" is now load-bearing and nothing asserts it. A rank-count allgather, or just a comment stating the invariant, would keep it from rotting.

Not approving: shared C++ runtime with no other approval yet, and 4 threads still open.

@liji-nv

liji-nv commented Jul 31, 2026

Copy link
Copy Markdown
Collaborator Author

/bot run --disable-fail-fast

@tensorrt-cicd

Copy link
Copy Markdown
Collaborator

PR_Github #62935 [ run ] triggered by Bot. Commit: 3ec5cbf Link to invocation

@BowenFu

BowenFu commented Jul 31, 2026

Copy link
Copy Markdown

Re-reviewed 3ec5cbfef. The idempotency guard and the combined deallocation+manager error are both good. One follow-up, same class as the residual I raised on 427b5823c — not blocking.

shutdown_userbuffers() now documents its contract explicitly (model_engine.py:2627-2632: "must be called by a deterministic executor shutdown path on every rank"), and unregister_user_buffer_collective barriers inside it. But there are now three rank-local early exits in front of that collective:

  1. model_engine.py:2634_userbuffers_shutdown_failed raises before entering.
  2. model_engine.py:2642not ub_buffers and not manager_initialized returns before entering.
  3. encoder_executor.py — the new _cleanup_done returns before entering.

Each is correct only if its predicate is rank-uniform. Where it isn't, the failure mode is a hang, not an error: the skipping rank exits while its peers sit in the barrier. (1) is the sharpest — a per-rank deallocation failure is exactly the case where ranks diverge, and that's the path that raises without the barrier.

I don't think any of these is reachable today (_init_userbuffers is config-driven and uniform across TP, and shutdown is called once per rank), so this is a hardening ask, not a bug report. But the collective contract is now load-bearing and nothing asserts it. Cheapest option is a comment on each guard stating the rank-uniformity assumption; if you want it enforced, an allreduce-of-the-predicate before the barrier would turn a hang into an error.

Signed-off-by: Jin Li <59594262+liji-nv@users.noreply.github.com>
@tensorrt-cicd

Copy link
Copy Markdown
Collaborator

PR_Github #62935 [ run ] completed with state FAILURE. Commit: 3ec5cbf
/LLM/main/L0_MergeRequest_PR pipeline #51055 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

@liji-nv

liji-nv commented Jul 31, 2026

Copy link
Copy Markdown
Collaborator Author

/bot run --disable-fail-fast

@tensorrt-cicd

Copy link
Copy Markdown
Collaborator

PR_Github #63036 [ run ] triggered by Bot. Commit: b2b1597 Link to invocation

@tensorrt-cicd

Copy link
Copy Markdown
Collaborator

PR_Github #63036 [ run ] completed with state FAILURE. Commit: b2b1597
/LLM/main/L0_MergeRequest_PR pipeline #51142 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

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants