From c77f0384704be19676f33b2bd3b4bde57cc956ed Mon Sep 17 00:00:00 2001 From: handongl Date: Mon, 27 Jul 2026 10:44:07 -0700 Subject: [PATCH 1/7] [nvbugs/6501404][fix] Gate NCCL symmetric output window on the registration threshold runNCCLAllReduceSymmetric allocated its window-backed output buffer with an unconditional createNCCLWindowTensor, bypassing the minRegistrationThreshold gate that the input path a few lines above already applies. That threshold is set to SIZE_MAX when neither NVLink nor MNNVL is supported, so on such topologies the collective ncclAllReduce plus cudaStreamSynchronize inside allocateAndRegisterBuffer never completes and every rank hangs, which the CI stage reports as "Test terminated unexpectedly". Apply the same gate to the output allocation: request a window buffer only when the input already obtained one or the message is at least as large as the threshold. Gating on the threshold rather than on mIsNVLINKSupported / mIsMNNVLSupported keeps the symmetric-memory fast path enabled wherever it works and still honors TLLM_NCCL_MIN_REGISTRATION. The pre-existing invalid-buffer fallback to torch::empty_like covers the skipped case, so no new error path is introduced. Signed-off-by: handongl --- cpp/tensorrt_llm/thop/allreduceOp.cpp | 23 +++++++++++++++++------ 1 file changed, 17 insertions(+), 6 deletions(-) diff --git a/cpp/tensorrt_llm/thop/allreduceOp.cpp b/cpp/tensorrt_llm/thop/allreduceOp.cpp index 85fbd59793d5..0035e30725ba 100644 --- a/cpp/tensorrt_llm/thop/allreduceOp.cpp +++ b/cpp/tensorrt_llm/thop/allreduceOp.cpp @@ -569,19 +569,30 @@ 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()) + // Use a window-backed output buffer under the same threshold gate as the input above. + // minRegistrationThreshold is SIZE_MAX without NVLink/MNNVL, where the collective + // ncclAllReduce plus cudaStreamSynchronize inside allocateAndRegisterBuffer cannot + // complete, so allocating here unconditionally hangs every rank. + torch::Tensor outputTensor; + if (windowBuffer0.isValid() || bufferSizeBytes >= minRegistrationThreshold) + { + 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)); // This operation owns only the temporary used to copy an unregistered input. The collective // is its final consumer, and same-stream release makes it reusable without synchronizing. From 08a87731dbe481d8be3b48a398ff47ee0650099b Mon Sep 17 00:00:00 2001 From: linquanh Date: Wed, 12 Aug 2026 06:29:47 +0000 Subject: [PATCH 2/7] [nvbugs/6501404][test] Move test_row_linear_norm_fusion to post_merge test_row_linear_norm_fusion needs 2 GPUs and has been flaky under nvbugs/6501404. Mark it @pytest.mark.post_merge so it is dropped from the blocking pre_merge 2-GPU glob (unittest/_torch/multi_gpu -m "not post_merge") on l0_dgx_h100, and add a matching 2-GPU post_merge condition block that collects it via -m "post_merge" so coverage is retained in the non-blocking post_merge stage. Signed-off-by: linquanh --- .../test_lists/test-db/l0_dgx_h100.yml | 16 ++++++++++++++++ tests/unittest/_torch/multi_gpu/test_linear.py | 1 + 2 files changed, 17 insertions(+) diff --git a/tests/integration/test_lists/test-db/l0_dgx_h100.yml b/tests/integration/test_lists/test-db/l0_dgx_h100.yml index 8354a7290642..4500d4331952 100644 --- a/tests/integration/test_lists/test-db/l0_dgx_h100.yml +++ b/tests/integration/test_lists/test-db/l0_dgx_h100.yml @@ -340,3 +340,19 @@ l0_dgx_h100: - unittest/_torch/moe/test_moe_module.py::test_configurable_moe_multi_gpu -k "CUTLASS and FP8 and not FP8_BLOCK_SCALES" - accuracy/test_llm_api_pytorch.py::TestNemotronV3Super::test_fp8_4gpus[attention_dp_off] - accuracy/test_llm_api_pytorch.py::TestNemotronV3Super::test_fp8_4gpus[attention_dp_on] +- condition: + 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) diff --git a/tests/unittest/_torch/multi_gpu/test_linear.py b/tests/unittest/_torch/multi_gpu/test_linear.py index d78a2998c579..fb71c4a2df24 100644 --- a/tests/unittest/_torch/multi_gpu/test_linear.py +++ b/tests/unittest/_torch/multi_gpu/test_linear.py @@ -322,6 +322,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}") From 44cacdb6de07e108773e4852a2a7a5a0d5ccac29 Mon Sep 17 00:00:00 2001 From: linquanh Date: Thu, 13 Aug 2026 03:08:09 +0000 Subject: [PATCH 3/7] [nvbugs/6501404][fix] Gate output window on uniform threshold Signed-off-by: linquanh --- cpp/tensorrt_llm/thop/allreduceOp.cpp | 9 ++++----- 1 file changed, 4 insertions(+), 5 deletions(-) diff --git a/cpp/tensorrt_llm/thop/allreduceOp.cpp b/cpp/tensorrt_llm/thop/allreduceOp.cpp index 0035e30725ba..6aaeb5fdd520 100644 --- a/cpp/tensorrt_llm/thop/allreduceOp.cpp +++ b/cpp/tensorrt_llm/thop/allreduceOp.cpp @@ -569,12 +569,11 @@ class AllreduceOp inputPtr = windowBuffer0.ptr; } - // Use a window-backed output buffer under the same threshold gate as the input above. - // minRegistrationThreshold is SIZE_MAX without NVLink/MNNVL, where the collective - // ncclAllReduce plus cudaStreamSynchronize inside allocateAndRegisterBuffer cannot - // complete, so allocating here unconditionally hangs every rank. + // 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 (windowBuffer0.isValid() || bufferSizeBytes >= minRegistrationThreshold) + if (bufferSizeBytes >= minRegistrationThreshold) { auto [windowOutput, windowBuffer1] = createNCCLWindowTensor(rawComm, input.sizes(), input.scalar_type()); if (windowBuffer1.isValid()) From eae1533cb2e3eafcce4ff9cbf69e69153da4e816 Mon Sep 17 00:00:00 2001 From: linquanh Date: Thu, 13 Aug 2026 07:56:14 +0000 Subject: [PATCH 4/7] [nvbugs/6501404][test] Add H100 post-merge stage Signed-off-by: linquanh --- jenkins/L0_Test.groovy | 1 + 1 file changed, 1 insertion(+) diff --git a/jenkins/L0_Test.groovy b/jenkins/L0_Test.groovy index df08274ef5df..8ac2f45b3242 100644 --- a/jenkins/L0_Test.groovy +++ b/jenkins/L0_Test.groovy @@ -6562,6 +6562,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], From 48322c6f95d1648922a4172e373837e1c7e8b1be Mon Sep 17 00:00:00 2001 From: handongl Date: Mon, 27 Jul 2026 13:00:55 -0700 Subject: [PATCH 5/7] [nvbugs/6501404][chore] Remove stale waiver after fix Signed-off-by: handongl [nvbugs/6501404][chore] Sort ray_orchestrator waive entry The file-contents-sorter pre-commit hook requires waives.txt to be byte-sorted. The newly added ray_orchestrator waive was placed after the sampler entry; reorder it before to satisfy the hook. Signed-off-by: linquanh --- tests/integration/test_lists/waives.txt | 1 - 1 file changed, 1 deletion(-) diff --git a/tests/integration/test_lists/waives.txt b/tests/integration/test_lists/waives.txt index 2dbc48f0c8da..baa757f613b8 100644 --- a/tests/integration/test_lists/waives.txt +++ b/tests/integration/test_lists/waives.txt @@ -185,7 +185,6 @@ unittest/_torch/modules/mamba/test_flashinfer_gdn_decode.py SKIP (https://nvbugs unittest/_torch/modules/tests_lora_modules/test_nemotron_h_lora_sanity.py::TestNemotronHLoRA::test_lora_pp2_sanity SKIP (https://nvbugs/6428124) 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/moe/test_moe_backend.py::test_trtllm_bf16_unquantized_moe[deepseekv3-relu2-8-separated_routing] SKIP (https://nvbugs/6778221) -unittest/_torch/multi_gpu/test_linear.py::test_row_linear_norm_fusion[2-hidden:16-seqlen:2] SKIP (https://nvbugs/6501404) unittest/_torch/ray_orchestrator/multi_gpu/test_accuracy_with_allreduce_strategy.py::test_accuracy_with_allreduce_strategy[AUTO-Qwen3/Qwen3-8B] SKIP (https://nvbugs/6923016) unittest/_torch/ray_orchestrator/multi_gpu/test_llm_update_weights_multi_gpu.py::test_llm_update_weights_qwen35_35b_bf16 SKIP (https://nvbugs/6907190) unittest/_torch/test_connector_adp.py SKIP (https://nvbugs/7003528) From 4c6312d8605dc3934afbaa36fd30d6e8e178b609 Mon Sep 17 00:00:00 2001 From: trtllm-agent <296075020+trtllm-agent@users.noreply.github.com> Date: Sat, 24 Oct 2026 06:13:33 -0700 Subject: [PATCH 6/7] [nvbugs/6501404][fix] Keep the output window handle in scope for the RMS_NORM release Rebasing onto main brings in the RMS_NORM release of the internal output window in runNCCLAllReduceSymmetric (#20232), which reads windowBuffer1 after the fused fallback. The threshold-gated output allocation declared windowBuffer1 inside the gate, so that read was out of scope. Declare windowBuffer1 at function scope. It stays invalid unless the gate allocates a valid window, which is exactly when outputTensor is window-backed and must be released, so the release keeps its meaning. Signed-off-by: trtllm-agent <296075020+trtllm-agent@users.noreply.github.com> --- cpp/tensorrt_llm/thop/allreduceOp.cpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/cpp/tensorrt_llm/thop/allreduceOp.cpp b/cpp/tensorrt_llm/thop/allreduceOp.cpp index 6aaeb5fdd520..3bc9d2979489 100644 --- a/cpp/tensorrt_llm/thop/allreduceOp.cpp +++ b/cpp/tensorrt_llm/thop/allreduceOp.cpp @@ -573,9 +573,11 @@ class AllreduceOp // registration threshold used by the input path. Smaller messages use a // regular output tensor and skip the window allocation path. torch::Tensor outputTensor; + tensorrt_llm::common::nccl_util::NCCLWindowBuffer windowBuffer1; if (bufferSizeBytes >= minRegistrationThreshold) { - auto [windowOutput, windowBuffer1] = createNCCLWindowTensor(rawComm, input.sizes(), input.scalar_type()); + torch::Tensor windowOutput; + std::tie(windowOutput, windowBuffer1) = createNCCLWindowTensor(rawComm, input.sizes(), input.scalar_type()); if (windowBuffer1.isValid()) { outputTensor = windowOutput; From 7fc43aa6da28000f1e0b5bf291f820ceed9790c8 Mon Sep 17 00:00:00 2001 From: trtllm-agent <296075020+trtllm-agent@users.noreply.github.com> Date: Wed, 28 Oct 2026 02:13:02 -0700 Subject: [PATCH 7/7] [nvbugs/6501404][chore] Revert all changes except the waiver removal The failure does not reproduce, so this PR now only removes the test_row_linear_norm_fusion[2-hidden:16-seqlen:2] waiver. Revert the other changes it carried: - the output-window gating in allreduceOp.cpp ("Gate NCCL symmetric output window on the registration threshold", "Gate output window on uniform threshold" and "Keep the output window handle in scope for the RMS_NORM release"); - the @pytest.mark.post_merge marker on test_row_linear_norm_fusion and the 2-GPU H100 post_merge block in l0_dgx_h100.yml ("Move test_row_linear_norm_fusion to post_merge"); - the DGX_H100-2_GPUs-PyTorch-Others-Post-Merge-1 stage in L0_Test.groovy ("Add H100 post-merge stage"). allreduceOp.cpp, test_linear.py, l0_dgx_h100.yml and L0_Test.groovy now match the base branch, so test_row_linear_norm_fusion runs again in the existing pre-merge `unittest/_torch/multi_gpu -m "not post_merge"` entry. Signed-off-by: trtllm-agent <296075020+trtllm-agent@users.noreply.github.com> --- cpp/tensorrt_llm/thop/allreduceOp.cpp | 24 +++++-------------- jenkins/L0_Test.groovy | 1 - .../test_lists/test-db/l0_dgx_h100.yml | 16 ------------- .../unittest/_torch/multi_gpu/test_linear.py | 1 - 4 files changed, 6 insertions(+), 36 deletions(-) diff --git a/cpp/tensorrt_llm/thop/allreduceOp.cpp b/cpp/tensorrt_llm/thop/allreduceOp.cpp index 3bc9d2979489..85fbd59793d5 100644 --- a/cpp/tensorrt_llm/thop/allreduceOp.cpp +++ b/cpp/tensorrt_llm/thop/allreduceOp.cpp @@ -569,31 +569,19 @@ class AllreduceOp inputPtr = windowBuffer0.ptr; } - // 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; - tensorrt_llm::common::nccl_util::NCCLWindowBuffer windowBuffer1; - if (bufferSizeBytes >= minRegistrationThreshold) - { - torch::Tensor windowOutput; - std::tie(windowOutput, windowBuffer1) = createNCCLWindowTensor(rawComm, input.sizes(), input.scalar_type()); - if (windowBuffer1.isValid()) - { - outputTensor = windowOutput; - } - } - if (!outputTensor.defined()) + // 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()) { 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, outputTensor.data_ptr(), size, (*getDtypeMap())[mType], ncclSum, comm, stream)); + NCCLCHECK_THROW(ncclAllReduce(inputPtr, outputPtr, size, (*getDtypeMap())[mType], ncclSum, comm, stream)); // This operation owns only the temporary used to copy an unregistered input. The collective // is its final consumer, and same-stream release makes it reusable without synchronizing. diff --git a/jenkins/L0_Test.groovy b/jenkins/L0_Test.groovy index 8ac2f45b3242..df08274ef5df 100644 --- a/jenkins/L0_Test.groovy +++ b/jenkins/L0_Test.groovy @@ -6562,7 +6562,6 @@ 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], diff --git a/tests/integration/test_lists/test-db/l0_dgx_h100.yml b/tests/integration/test_lists/test-db/l0_dgx_h100.yml index 4500d4331952..8354a7290642 100644 --- a/tests/integration/test_lists/test-db/l0_dgx_h100.yml +++ b/tests/integration/test_lists/test-db/l0_dgx_h100.yml @@ -340,19 +340,3 @@ l0_dgx_h100: - unittest/_torch/moe/test_moe_module.py::test_configurable_moe_multi_gpu -k "CUTLASS and FP8 and not FP8_BLOCK_SCALES" - accuracy/test_llm_api_pytorch.py::TestNemotronV3Super::test_fp8_4gpus[attention_dp_off] - accuracy/test_llm_api_pytorch.py::TestNemotronV3Super::test_fp8_4gpus[attention_dp_on] -- condition: - 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) diff --git a/tests/unittest/_torch/multi_gpu/test_linear.py b/tests/unittest/_torch/multi_gpu/test_linear.py index fb71c4a2df24..d78a2998c579 100644 --- a/tests/unittest/_torch/multi_gpu/test_linear.py +++ b/tests/unittest/_torch/multi_gpu/test_linear.py @@ -322,7 +322,6 @@ 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}")