Skip to content

[https://nvbugs/6501404][fix] Request the output window only when bufferSizeBytes >= minRegistrationThreshold #16911

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
trtllm-agent wants to merge 5 commits into
base: main
Choose a base branch
from
Open

[https://nvbugs/6501404][fix] Request the output window only when bufferSizeBytes >= minRegistrationThreshold #16911

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
b4644bd
[nvbugs/6501404][fix] Gate NCCL symmetric output window on the regist…
HandongLi-01 Jul 27, 2026
0945dda
[nvbugs/6501404][test] Move test_row_linear_norm_fusion to post_merge
kris1025 Aug 12, 2026
b98caf0
[nvbugs/6501404][fix] Gate output window on uniform threshold
kris1025 Aug 13, 2026
7a12a7c
[nvbugs/6501404][test] Add H100 post-merge stage
kris1025 Aug 13, 2026
a7161f4
[nvbugs/6501404][chore] Remove stale waiver after fix
HandongLi-01 Jul 27, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
22 changes: 16 additions & 6 deletions cpp/tensorrt_llm/thop/allreduceOp.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -563,19 +563,29 @@ class AllreduceOp
inputPtr = windowBuffer0.ptr;
}

// Use window-backed output buffer
auto [normOut, windowBuffer1] = createNCCLWindowTensor(rawComm, input.sizes(), input.scalar_type());
torch::Tensor outputTensor = windowBuffer1.isValid() ? normOut : torch::empty_like(inputTensor);
void* outputPtr = windowBuffer1.isValid() ? windowBuffer1.ptr : outputTensor.data_ptr();
if (!windowBuffer1.isValid())
// Request an output window only when the message meets the same
// registration threshold used by the input path. Smaller messages use a
// regular output tensor and skip the window allocation path.
torch::Tensor outputTensor;
if (bufferSizeBytes >= minRegistrationThreshold)

@crazydemo crazydemo Aug 26, 2026

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.

[MINOR] Output gate drops the windowBuffer0.isValid() branch used by input

The output window is now requested only when bufferSizeBytes >= minRegistrationThreshold. The release notes describe the intended condition as windowBuffer0.isValid() || bufferSizeBytes >= minRegistrationThreshold, matching the input path. As written, if the input is window-backed (windowBuffer0 valid) but the message is below the threshold, the input uses a registered symmetric pointer while the output falls back to a plain torch::empty_like pointer. Mixing a registered send buffer with an unregistered recv buffer in runNCCLAllReduceSymmetric is exactly the asymmetry class this code is sensitive to and could degrade or misbehave on the symmetric kernel. Confirm whether the input path can be window-backed for sub-threshold messages; if so, mirror the full condition here.

@kris1025 kris1025 Aug 26, 2026

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.

The description is updated now.

{
auto [windowOutput, windowBuffer1] = createNCCLWindowTensor(rawComm, input.sizes(), input.scalar_type());
if (windowBuffer1.isValid())
{
outputTensor = windowOutput;
}
}
if (!outputTensor.defined())
{
TLLM_LOG_DEBUG(
"[runNCCLAllReduceSymmetric] No valid symmetric buffer available; "
"using plain CUDA tensor for output");
outputTensor = torch::empty_like(inputTensor);
}

// Perform allreduce
NCCLCHECK_THROW(ncclAllReduce(inputPtr, outputPtr, size, (*getDtypeMap())[mType], ncclSum, comm, stream));
NCCLCHECK_THROW(
ncclAllReduce(inputPtr, outputTensor.data_ptr(), size, (*getDtypeMap())[mType], ncclSum, comm, stream));
Comment thread
kris1025 marked this conversation as resolved.

if (mOp == AllReduceFusionOp::NONE)
{
Expand Down
1 change: 1 addition & 0 deletions jenkins/L0_Test.groovy
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6193,6 +6193,7 @@ def launchTestJobs(pipeline, testFilter, globalVars)
"DGX_A100-FMHA-Post-Merge-1": ["auto:dgx-a100-x1", "l0_a100", 1, 1],
"DGX_H100-2_GPUs-PyTorch-Others-1": ["auto:dgx-h100-x2", "l0_dgx_h100", 1, 2, 2],
"DGX_H100-2_GPUs-PyTorch-Others-2": ["auto:dgx-h100-x2", "l0_dgx_h100", 2, 2, 2],
"DGX_H100-2_GPUs-PyTorch-Others-Post-Merge-1": ["auto:dgx-h100-x2", "l0_dgx_h100", 1, 1, 2],
"DGX_H100-2_GPUs-PyTorch-GptOss-1": ["auto:dgx-h100-x2", "l0_dgx_h100", 1, 1, 2],
"DGX_H100-2_GPUs-PyTorch-Ray-1": ["auto:dgx-h100-x2", "l0_dgx_h100", 1, 1, 2],
"DGX_H100-4_GPUs-PyTorch-DeepSeek-1": ["auto:dgx-h100-x4", "l0_dgx_h100", 1, 1, 4],
Expand Down
16 changes: 16 additions & 0 deletions tests/integration/test_lists/test-db/l0_dgx_h100.yml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -343,3 +343,19 @@ l0_dgx_h100:
- accuracy/test_llm_api_pytorch.py::TestNemotronV3Super::test_fp8_4gpus[attention_dp_off-cpp_mamba_cache]
- accuracy/test_llm_api_pytorch.py::TestNemotronV3Super::test_fp8_4gpus[attention_dp_on-python_mamba_cache]
- accuracy/test_llm_api_pytorch.py::TestNemotronV3Super::test_fp8_4gpus[attention_dp_on-cpp_mamba_cache]
- condition:

@chzblych chzblych Aug 13, 2026

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.

We don't have a test stage matching those configs. You need to add a new test stage like DGX_H100-2_GPUs-PyTorch-Others-Post-Merge-1 here - https://github.com/NVIDIA/TensorRT-LLM/blob/main/jenkins/L0_Test.groovy#L5443.

@kris1025 kris1025 Aug 13, 2026

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.

The new stage is added now.

ranges:
system_gpu_count:
gte: 2
lte: 2
wildcards:
gpu:
- '*h100*'
linux_distribution_name: ubuntu*
terms:
stage: post_merge
backend: pytorch
auto_trigger: others
orchestrator: mpi
tests:
- unittest/_torch/multi_gpu -m "post_merge" TIMEOUT (90)
1 change: 0 additions & 1 deletion tests/integration/test_lists/waives.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -345,7 +345,6 @@ unittest/_torch/modeling/test_modeling_nemotron_nano_v2_vl.py::test_nemotron_nan
unittest/_torch/modules/tests_lora_modules/test_nemotron_h_lora_sanity.py::TestNemotronHLoRA::test_lora_pp2_sanity SKIP (https://nvbugs/6428124)
unittest/_torch/modules/tests_lora_modules/test_qwen3_sanity.py::TestQwen3LoRA::test_qwen3_fp8_lora SKIP (https://nvbugs/6668777)
unittest/_torch/moe/test_moe_backend.py::test_moe_backend[act=Relu2-e60_k4_h2048_i1408-seq=8-dtype=torch.bfloat16-backend=TRTLLM-quant=NVFP4-routing=Renormalize] SKIP (https://nvbugs/5989912)
unittest/_torch/multi_gpu/test_linear.py::test_row_linear_norm_fusion[2-hidden:16-seqlen:2] SKIP (https://nvbugs/6501404)
unittest/_torch/speculative/hw_agnostic/test_dflash.py::test_dflash_qwen3_5_4b[False] SKIP (https://nvbugs/6535767)
unittest/_torch/speculative/hw_agnostic/test_dflash.py::test_dflash_qwen3_5_4b[True] SKIP (https://nvbugs/6535767)
unittest/_torch/speculative/hw_agnostic/test_ngram.py::test_llama_ngram[True-True-TRTLLM] SKIP (https://nvbugs/6507102)
Expand Down
1 change: 1 addition & 0 deletions tests/unittest/_torch/multi_gpu/test_linear.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -310,6 +310,7 @@ def test_row_linear(hidden_size, mpi_pool_executor):
assert r is True


@pytest.mark.post_merge
@pytest.mark.skipif(torch.cuda.device_count() < 2,
reason='needs 2 GPUs to run this test')
@pytest.mark.parametrize("seq_len", [2, 32], ids=lambda x: f"seqlen:{x}")
Expand Down
Loading