From b4644bd33c68170509ac55734ed96053d18695e0 Mon Sep 17 00:00:00 2001 From: handongl Date: Mon, 27 Jul 2026 10:44:07 -0700 Subject: [PATCH 1/5] [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 fe95722ab5dc..2b08591c70ea 100644 --- a/cpp/tensorrt_llm/thop/allreduceOp.cpp +++ b/cpp/tensorrt_llm/thop/allreduceOp.cpp @@ -563,19 +563,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)); if (mOp == AllReduceFusionOp::NONE) { From 0945dda5fac7f684b05ab916d11805133329d798 Mon Sep 17 00:00:00 2001 From: linquanh Date: Wed, 12 Aug 2026 06:29:47 +0000 Subject: [PATCH 2/5] [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 43ab57a68f7a..83997d5341a5 100644 --- a/tests/integration/test_lists/test-db/l0_dgx_h100.yml +++ b/tests/integration/test_lists/test-db/l0_dgx_h100.yml @@ -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: + 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 4db1c6b053ce..3e4b40b2f57c 100644 --- a/tests/unittest/_torch/multi_gpu/test_linear.py +++ b/tests/unittest/_torch/multi_gpu/test_linear.py @@ -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}") From b98caf09105213fdb98af7ebf4457c181aaefe16 Mon Sep 17 00:00:00 2001 From: linquanh Date: Thu, 13 Aug 2026 03:08:09 +0000 Subject: [PATCH 3/5] [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 2b08591c70ea..06332fc50407 100644 --- a/cpp/tensorrt_llm/thop/allreduceOp.cpp +++ b/cpp/tensorrt_llm/thop/allreduceOp.cpp @@ -563,12 +563,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 7a12a7c96018bbadd86dd05a140a6d7e9617d2b9 Mon Sep 17 00:00:00 2001 From: linquanh Date: Thu, 13 Aug 2026 07:56:14 +0000 Subject: [PATCH 4/5] [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 3911a2e9151d..37c5e2efe419 100644 --- a/jenkins/L0_Test.groovy +++ b/jenkins/L0_Test.groovy @@ -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], From a7161f44554a218ccaab9c3de48480cf3082a25a Mon Sep 17 00:00:00 2001 From: handongl Date: Mon, 27 Jul 2026 13:00:55 -0700 Subject: [PATCH 5/5] [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 c3a649190f0f..72b5d5893b38 100644 --- a/tests/integration/test_lists/waives.txt +++ b/tests/integration/test_lists/waives.txt @@ -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)