From 86ce503474f1ccf73d1610fe69a0c7a97ee7e10f Mon Sep 17 00:00:00 2001 From: Shiyu Li Date: Wed, 22 Oct 2025 23:07:03 -0700 Subject: [PATCH 1/2] Fix hopper accuracy issue with wrong sem Signed-off-by: Shiyu Li --- .../communicationKernels/mnnvlTwoShotAllreduceKernels.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/tensorrt_llm/kernels/communicationKernels/mnnvlTwoShotAllreduceKernels.cu b/cpp/tensorrt_llm/kernels/communicationKernels/mnnvlTwoShotAllreduceKernels.cu index c38abd957850..f439f2b8efd0 100644 --- a/cpp/tensorrt_llm/kernels/communicationKernels/mnnvlTwoShotAllreduceKernels.cu +++ b/cpp/tensorrt_llm/kernels/communicationKernels/mnnvlTwoShotAllreduceKernels.cu @@ -117,8 +117,8 @@ __device__ struct __attribute__((aligned(32))) LamportFlags { #if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 1000)) asm volatile("red.async.release.global.gpu.add.u32 [%0], %1;" ::"l"(offset_access_ptr), "r"(1) : "memory"); -#elif (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)) - asm volatile("red.global.gpu.add.u32 [%0], %1;" ::"l"(offset_access_ptr), "r"(1) : "memory"); +#elif (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 700)) + asm volatile("red.release.global.gpu.add.u32 [%0], %1;" ::"l"(offset_access_ptr), "r"(1) : "memory"); #else atomicAdd(offset_access_ptr, 1); #endif From f8b1fb051bfb29c3ef979cadf473b533d4f30ed3 Mon Sep 17 00:00:00 2001 From: Shiyu Li Date: Wed, 22 Oct 2025 23:13:20 -0700 Subject: [PATCH 2/2] Add back the unit test Signed-off-by: Shiyu Li --- tests/unittest/_torch/multi_gpu/test_mnnvl_allreduce.py | 1 - 1 file changed, 1 deletion(-) diff --git a/tests/unittest/_torch/multi_gpu/test_mnnvl_allreduce.py b/tests/unittest/_torch/multi_gpu/test_mnnvl_allreduce.py index 7a3560cc14b7..53692c519216 100644 --- a/tests/unittest/_torch/multi_gpu/test_mnnvl_allreduce.py +++ b/tests/unittest/_torch/multi_gpu/test_mnnvl_allreduce.py @@ -183,7 +183,6 @@ def func(input, residual, norm_weight, eps, enable_fusion): ) def test_row_linear_residual_norm_fusion(seq_len, hidden_size, dtype, strategy, fusion): - pytest.skip("https://nvbugs/5597647") torch.manual_seed(42) tensor_parallel_size = 2