Skip to content

feat(trtllm): Rubin build support and rollout fixes - #4123

Draft
hchings wants to merge 32 commits into
mainfrom
erinh/trtllm-mlperf-mislan
Draft

hchings wants to merge 32 commits into
mainfrom
erinh/trtllm-mlperf-mislan

Conversation

@hchings

@hchings hchings commented Sep 14, 2026

Copy link
Copy Markdown
Contributor

What does this PR do ?

Add a one line overview of what this PR aims to accomplish.

Issues

List issues that this PR closes (syntax):

Usage

  • You can potentially add a usage example below
# Add a code snippet demonstrating how to use this

Before your PR is "Ready for review"

Pre checks:

  • Make sure you read and followed Contributor guidelines
  • Did you write any new necessary tests?
  • Did you run the unit tests and functional tests locally? Visit our Testing Guide for how to run tests
  • Did you add or update any necessary documentation? Visit our Document Development Guide for how to write, build and test the docs.

Additional Information

  • ...

shuyixiong and others added 30 commits September 13, 2026 17:10
Signed-off-by: shuyixiong <219646547+shuyixiong@users.noreply.github.com>
Signed-off-by: Erin Ho <erinh@nvidia.com>
Signed-off-by: Erin Ho <14718778+hchings@users.noreply.github.com>
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
Co-authored-by: Erin Ho <erinh@nvidia.com>
Co-authored-by: Erin Ho <14718778+hchings@users.noreply.github.com>
#3428)

Signed-off-by: Superjomn <yanchunwei@outlook.com>
Co-authored-by: Erin Ho <14718778+hchings@users.noreply.github.com>
The HTTP server built its TrtSamplingParams from temperature and top_p
only, so a recipe setting policy.generation.top_k was silently ignored
whenever rollouts went through the server -- which is every NeMo Gym and
PD-disaggregated run. The direct path (_build_sampling_params) has always
applied it, so the two paths sampled from different distributions for the
same config, and a run that switched between them was not comparable with
itself.

Pass top_k through, using the same convention as the direct path: TRT-LLM
spells "no top-k restriction" as 0 while the generation config spells it
as null. Add it to the request-validation loop too, so a request that
disagrees with the server's config is rejected rather than quietly
overridden -- the same treatment temperature and top_p already get. That
loop now reads sampling_config with .get(): top_k is absent from configs
written before this change, and a KeyError there would reject every
request instead of the mismatched ones.

Also set logprobs_simple_format=True. Without it TRT-LLM returns the
verbose logprob structure, which the adapter has to walk per token; the
simple format is what the direct path consumes.

Signed-off-by: shuyixiong <219646547+shuyixiong@users.noreply.github.com>
Signed-off-by: shuyixiong <219646547+shuyixiong@users.noreply.github.com>
…kens

The Rubin TRT-LLM base image ships PIP_CONSTRAINT=/etc/pip/constraint.txt,
pinning the versions vendored into the image's own Python. build_wheel.py's
setup_venv() shells out to real pip, which honours it, and those versions are
not published anywhere: the image pins cuda-python==13.4.0, which does not
exist on PyPI (latest 13.x is 13.3.1, already installed by uv). So the
otherwise-satisfied `cuda-python>=13` in requirements.txt became unresolvable
and the build died before cmake ran. Unset the constraint -- this venv is
uv-managed and owes the base image's site-packages nothing.

This only appeared after moving to the NGC-derived Rubin base; cuda-dl-base
does not set PIP_CONSTRAINT.

Also redact the clone token from two places that printed it verbatim into the
build log: the "TRT-LLM Git URL" echo, and _backend.py's CalledProcessError,
which stringifies the whole argv (the expanded url included) on failure.

And drop nvidia-modelopt from requirements.txt, preventively: build_wheel.py
pip-installs into the live nemo-rl venv, so the ref's `~=0.39.0` pin would let
pip downgrade the modelopt uv resolved from a git rev at 0.46.0.dev*.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: shuyixiong <219646547+shuyixiong@users.noreply.github.com>
Adding 10.7 to TORCH_CUDA_ARCH_LIST / HYBRID_EP_CUDA_ARCH_LIST fails to build:
the torch extensions compiled during `uv sync` (mamba-ssm, causal-conv1d,
transformer-engine) and DeepEP cannot target sm_107 with this toolchain. Fall
back to Blackwell, matching what the opt/dev-backup-rubin reference shipped.

TRT-LLM itself still targets Rubin via _DEFAULT_ARCH / BUILD_CUSTOM_TRTLLM_ARCH
=107-real; only the torch-side extensions are affected here.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: shuyixiong <219646547+shuyixiong@users.noreply.github.com>
Signed-off-by: shuyixiong <219646547+shuyixiong@users.noreply.github.com>
Extend the default SM arch list to 100-real;103-real;107-real (GB200/B200,
GB300/B300, Rubin) in both places that must agree: _DEFAULT_ARCH in _backend.py,
which is folded into the wheel cache key, and the ARCH fallback in
build-custom-trtllm.sh.

The nvshmem patch gets the same three archs as 100\;103\;107 -- bare, because
nvshmem rejects the suffixed names CMake generates, which is the reason that
patch exists. Its semicolons stay backslash-escaped: the string lands inside
CMAKE_CACHE_ARGS of an ExternalProject_Add, where an unescaped ';' would split
one cache entry into three arguments and silently build sm_100 only.

That patch does not read BUILD_CUSTOM_TRTLLM_ARCH, so it and _DEFAULT_ARCH have
to be edited together; both comments now say so.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: shuyixiong <219646547+shuyixiong@users.noreply.github.com>
The pinned ref's requirements.txt carries tensorrt~=10.16.1, but it was absent
from [project].dependencies, so uv never resolved it -- the wheel build merely
pip-installed it transiently into the venv via requirements-dev.txt. Declare it
so the version uv manages matches the ref and the base image's TRT 10.16.1.11.

tensorrt on PyPI ships only an sdist whose wheel_stub downloads multi-GB
binaries at build time, which would make `uv lock` fetch them. Supply static
[[tool.uv.dependency-metadata]] for tensorrt / tensorrt-cu13 / -libs /
-bindings so the graph resolves from metadata alone; the real binaries are
fetched at `uv sync --extra trtllm` time inside Docker. Same approach, same
placement and versions as origin/main -- our base d5fb8d0 predates it.

Lock goes 553 -> 557 packages (the four tensorrt entries); `uv lock --check`
passes with submodules at their committed pins.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: shuyixiong <219646547+shuyixiong@users.noreply.github.com>
Signed-off-by: shuyixiong <219646547+shuyixiong@users.noreply.github.com>
torch_dist checkpointing opens with a save-planning gather_object over the
default process group: Megatron's dist_checkpointing/strategies/torch.py passes
process_group=None into save_state_dict_async_plan, and DCP's _DistWrapper
forwards that to dist.gather_object. On a NCCL default group that is a device
collective, and NCCL allocates its buffers with its own cudaMalloc rather than
through PyTorch's caching allocator, so memory the allocator is merely holding
is unreachable to it. With the device full the cudaMalloc fails and NCCL
surfaces it as the opaque "NCCL Error 1: unhandled cuda error" -- no torch OOM
is raised anywhere, because torch never requested the memory.

Job 2724162 hit this at its first save: every training rank sat at 277.5 GiB of
277.5 GiB and the run died in that gather. Reclaiming first measured 68 GiB of
headroom per rank (device-free 9.8 -> 77.7 GiB) and the 740 GiB step_1
checkpoint then completed.

nccl_reshard_refit and prepare_for_lp_inference already clear the cache ahead of
their own large NCCL phases; save_checkpoint was the path that skipped it. The
[CKPT_MEM] line is kept because the failure it guards against is silent -- NCCL
reports only "unhandled cuda error", so a recurrence would otherwise leave
nothing to reason from.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Signed-off-by: shuyixiong <219646547+shuyixiong@users.noreply.github.com>
Signed-off-by: Erin Ho <14718778+hchings@users.noreply.github.com>
Ports 4a910e9 and d9c9d61 (Michal Futrega, branch trtllm-agentic-swe-mlperf)
onto opt/dev, which has neither.

The default 10-minute NCCL watchdog SIGABRTs ranks that wait on peers stuck in
transient first-hit stalls -- triton JIT of GDN kernels for new packed-sequence
shapes being the prime suspect -- observed as coordinated "Terminating the
process after attempting to dump debug info" during step-2 collectives in jobs
2274348/2280396. Step 1 always passed, so these are warm-up effects rather than
real hangs; true hangs stay bounded by the SLURM walltime.

Both call sites are needed. setup_distributed() raises the default group's
timeout, which sub-groups normally inherit, but initialize_megatron creates the
TP/PP/EP/DP sub-groups with an EXPLICIT dist.distributed_timeout_minutes (mcore
default 10) that overrides that inheritance -- job 2282178 still died at
Timeout(ms)=600000 with only the first fix. Both read the same
NRL_NCCL_TIMEOUT_MINUTES env var, default 60.

Applied by hand rather than cherry-picked: the original commits sit on a base
whose init_process_group already passed device_id=cuda:{local_rank}, which
opt/dev does not have. Only the timeout change is taken, so the device-binding
difference between the branches is left alone.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Signed-off-by: shuyixiong <219646547+shuyixiong@users.noreply.github.com>
Moves the custom TensorRT-LLM wheel to user/zongfeij/rl @ 7035705878,
which adds recompute_active_requests and the Qwen3.5 MXFP8 refit path.

Still 1.3.0rc23, so [project].version and the root pyproject's trtllm
extra pin are unchanged. [project].dependencies are left as-is: the
ref's requirements.txt relaxes the tensorrt, flashinfer-python,
nvidia-cutlass-dsl and apache-tvm-ffi pins and adds
nvidia-cuda-nvrtc==13.4.46rc1 (all from the 0.8 drop container update),
but re-syncing those is deferred to a separate change.

Signed-off-by: Shuyi Xiong <219646547+shuyixiong@users.noreply.github.com>
Launchers already export MOUNT_LOG_DIR_IN_CONTAINER=1 alongside BASE_LOG_DIR
(qwen35_397b_grpo_trtllm/launch_ray_cluster.sh does), but nothing here read it,
so /logs stayed container-local: a per-node tmpfs in the enroot data dir.

Anything written under /logs on one node was therefore invisible to every
other node. NeMo-Gym is the case that matters -- its server resolves agent
trajectories and llm_completions under /logs while the Ray runners that
produce them are spread across the allocation, so the server found nothing,
every rollout came back with zero output items, and the trajectory collector
failed the step. mlperf's run_and_time.sh and logger.log_dir hardcode /logs
the same way.

Mount LOG_DIR (already required to be on a shared filesystem -- the
STARTED_RAY_HEAD / ray_worker_units / ENDED signalling depends on it) at /logs
when the flag is set. Guarded on empty MOUNTS like the UV_CACHE_DIR_OVERRIDE
block above, and appended before COMMON_SRUN_ARGS consumes MOUNTS.

Side benefit: driver and per-agent apptainer logs now survive the job instead
of dying with the container's tmpfs.
Signed-off-by: Erin Ho <14718778+hchings@users.noreply.github.com>
Force fla-core (and its flash-linear-attention meta) to 0.5.2 via the
[tool.uv] override-dependencies. 0.5.2 carries the two Blackwell GDN
autotune restrictions that 0.5.1 lacks:

  - PR #953 (commit 3eeef4e): restrict the fwd h kernel
    (chunk_gated_delta_rule_fwd_kernel_h_blockdim64) to num_warps=2 on
    Blackwell — Triton tl.dot recurrence race (fla issue #945).
  - PR #1000 (commit ac6c648): restrict the gated-delta bwd autotune
    (prepare_wy_repr_bwd_kernel) on Blackwell.

megatron-core's dev extra pins flash-linear-attention==0.5.1, which is the
racy release; upstream 0.5.2 fixes it. The recipes keep FLA_TILELANG=0
(triton backend) — the tilelang GDN backward fault is not addressed by 0.5.2.
configure_generation_config gave vllm_cfg["load_format"] the value
"auto" if is_eval else "dummy" but left trtllm_cfg alone, so a TRT-LLM
engine read the full checkpoint at startup even though refit_policy_generation
overwrites every weight before the trajectory collector is allowed to issue
its first request. On a 397B engine that read is minutes of startup the run
throws away.

Mirror the vLLM contract: set trtllm_cfg["load_format"] in
configure_generation_config and pass it through to the AsyncLLM constructor.
precision="fp8" stays "dummy" in both modes -- the engine is built with the
quantized layout the BF16 refit populates, and configure_fp8_llm_kwargs
already rejects anything else -- so this changes nothing for FP8.

Evaluation keeps "auto": it has no refit to supply weights.
The guard pinned the exact private class name
`fastokens._compat._TokenizerShim`. fastokens 0.3.1 renamed it to
`fastokens._ConfiguredTokenizerShim`, so `patch_transformers()` succeeded but
the worker still refused to start:

  RuntimeError: NRL_USE_FASTOKENS=1, but the TRT-LLM HTTP tokenizer backend
  is 'fastokens._ConfiguredTokenizerShim'; expected
  'fastokens._compat._TokenizerShim'

What the check is for is whether the monkey-patch landed at all, so match on
the `fastokens.` module prefix instead of an internal class name that moves
with every release.
get_logprobs returned BatchedDataDict(...).to("cpu"), leaving the
device-to-host transfer to BatchedDataDict.to(); that path hung at scale.
Allocate an explicit pinned host tensor and copy into it before returning,
so the transfer is direct and the returned logprobs are pinned.

Applied from the optimized repo's carried patch
qwen35_397b_grpo_trtllm/pytorch/patches/nemo-rl/0010-fix-megatron-copy-logprobs-to-pinned-CPU-memory.patch
(optimized commit c885684bc), which ports 6c8c46cbc from qwen35_397b_grpo.

Signed-off-by: Shuyi Xiong <219646547+shuyixiong@users.noreply.github.com>
Rebases the RL tekit branch onto feat/rubin-bringup, which replaces the
feat_sm107 flashinfer source build with the 0.6.18 internal wheel and drops
the cutlass-dsl 4.5.0 pin (tekit ff4598531 / 559775ea8 / de0cf4d8a).

Two build-custom-trtllm.sh patches become no-ops on this ref and would have
aborted via assert_patch_target:
  - requirements.txt already ships `setuptools>=80`
  - cutlass_kernels/CMakeLists.txt dropped the `setup_library.py develop
    --user` execute_process in favour of PYTHONPATH for generate_kernels.py

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Picks up tekit 780b94e65, which passes the three trailing arguments
trtllm_paged_attention_decode gained in the base image's flashinfer
0.6.18+8c3bbc00. The image can now keep the base's flashinfer instead of
force-installing the older +cf3c3a3e build.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
CI passes GITLAB_CLONE_ACCESS_TOKEN already as "user:token", so the prefix
expanded to "oauth2:user:token". URL userinfo splits on the first colon, which
made the password "user:token" and every CI build failed the tekit fetch with
"HTTP Basic: Access denied. You must use a token instead of a password". A bare
token hid the bug locally, where "oauth2:glpat-x" is well formed.

The url now carries the token alone, matching how the optimized Dockerfile
passes it elsewhere; callers holding a bare token add their own username.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Moves [tool.trtllm].ref from 780b94e655 to 1c6e8f3dfb ("Leave generation-only
requests out of the post-refit KV recompute on disaggregated engines").

The other two pins that the file's own comment says move together stay put,
having been checked rather than assumed: tensorrt_llm/version.py reads
1.3.0rc23 at both refs, so `version` here and the `trtllm` extra in the root
pyproject.toml are already correct, and requirements.txt is unchanged between
the two refs, so [project].dependencies needs no re-sync.

This is not a fast-forward -- the refs have diverged, with 20+ commits on the
new side, mostly disaggregation work: in-flight weight updates on attention-DP
disagg engines, refit with piecewise CUDA graphs, the folded save-last prefill
that TLLM_MAMBA_FOLD_SAVE_LAST pins, and the FMHA trailing-argument fix. What
the old ref carried that the new one does not has not been audited.

The wheel cache is ref-scoped, so this forces a rebuild.

Signed-off-by: Shuyi Xiong <219646547+shuyixiong@users.noreply.github.com>
Picks up "[None][fix] Ship the end-of-prompt recurrent state to the generation
engine when the save-last fold is on" and everything else on
user/zongfeij/rl since 1c6e8f3dfb.
Signed-off-by: Shuyi Xiong <219646547+shuyixiong@users.noreply.github.com>
Moves [tool.trtllm].ref in 3rdparty/TensorRT-LLM-workspace/pyproject.toml from
854bcea144 to b1cecb7c9d. build_trtllm compiles the engine from this ref, so
the image picks the new TRT-LLM up on the next NEMO_RL_REVISION bump.

The sha appears only here; nothing else in the tree references it.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Signed-off-by: Shuyi Xiong <219646547+shuyixiong@users.noreply.github.com>
A Gym rollout can return without a single assistant turn (agent stalls
before its first completion and Gym's wall-clock timeout kills it).
NRL_SKIP_FAILING_EMPTY_ROLLOUT=1 stands the sample up as prompt-only and
masks it out of the loss instead of raising and losing the step's other
rollouts, surfaced as train/num_masked_seqs_by_empty_rollout. Default
"0" keeps the original raise.

This is general error-handling behavior that applies to every generation
backend, not something specific to prefill/decode disaggregation -- it
originated bundled into the disagg bring-up commit but was carved back
out (see PR #4095's cc12c14 revert) since it changes behavior across
the board. This commit restores it here, alongside the rest of the
non-fp8, non-disagg work.

Co-Authored-By: shuyixiong <219646547+shuyixiong@users.noreply.github.com>
Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
… omissions

This comment block was introduced alongside the disagg bring-up commit
and got dropped when that commit's pyproject.toml hunk was excluded from
this branch. It's pure documentation (no functional change) explaining
why PyNvVideoCodec and packaging aren't mirrored into [project.dependencies].

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
@hchings hchings self-assigned this Sep 14, 2026
@copy-pr-bot

copy-pr-bot Bot commented Sep 14, 2026

Copy link
Copy Markdown

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

tests/unit/models/generation/test_trtllm_backend.py (missing the trtllm/
subdir its siblings live in) was a stray copy of the two tests already in
tests/unit/models/generation/trtllm/test_trtllm_refit.py -- a rebase
artifact from splitting this branch's work into separate PRs. It also lacked
pytestmark = pytest.mark.trtllm, so any CI shard running
tests/unit/models/generation/ without the trtllm marker filter and without
tensorrt_llm installed would hit a collection-time ImportError, and its
patch.object(WorkerExtension, "finalize_weight_update", None) omitted
create=True, which raises AttributeError on a TRT-LLM build where that
attribute doesn't exist at all.

Signed-off-by: Erin Ho <14718778+hchings@users.noreply.github.com>
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.

3 participants