f(rccl/gin): fix inverted system-scope fence guard on HIP backends

projects/rccl/src/include/nccl_device/gin/anvil_sdma/gin_anvil_sdma.h

@@ -183,8 +183,10 @@ struct ncclGinApi_Put<NCCL_NET_DEVICE_GIN_ANVIL_SDMA> {
     if (!anvilCtxValid(rsCtx)) return;
     const int blockId = blockIdx.x + blockIdx.y * gridDim.x;
 
-    if ((required == cuda::thread_scope_system) && (given > required)) {
-      __threadfence_system();
+    // HIP thread_scope (hip_compat.h): system is the MAX value, so a caller that
+    // only guaranteed a weaker scope has given < required -> add a system fence.
+    if ((required == cuda::thread_scope_system) && (given < required)) {
+      NCCL_GIN_THREADFENCE_SYSTEM();
     }
 
     size_t threshold = loadConst(&rsCtx->sdmaThreshold);
@@ -293,8 +295,10 @@ struct ncclGinApi_PutValue<NCCL_NET_DEVICE_GIN_ANVIL_SDMA> {
     ncclGinAnvilSdmaMemHandle* dstMh = (ncclGinAnvilSdmaMemHandle*)dstWin;
     T tmp = srcVal;
 
-    if ((required == cuda::thread_scope_system) && (given > required)) {
-      __threadfence_system();
+    // HIP thread_scope (hip_compat.h): system is the MAX value, so a caller that
+    // only guaranteed a weaker scope has given < required -> add a system fence.
+    if ((required == cuda::thread_scope_system) && (given < required)) {
+      NCCL_GIN_THREADFENCE_SYSTEM();
     }
 
     size_t threshold = loadConst(&rsCtx->sdmaThreshold);

projects/rccl/src/include/nccl_device/gin/gdaki/gin_gdaki.h

@@ -89,8 +89,13 @@ NCCL_DEVICE_INLINE static void putImplMode(
       counter_laddr.key = loadConst(&gdaki->sink_buffer_lkey);
     }
 
-    // cuda::thread_scope_system has the lowest value
-    // DOCA guarantees SCOPE_GPU. Only add another release if SCOPE_SYSTEM is required.
+    // gdaki is CUDA-only (NCCL_GIN_GDAKI_ENABLE=0 on HIP). In libcu++,
+    // cuda::thread_scope_system is the LOWEST value -- the opposite of HIP's
+    // hip_compat.h, where it is the max. DOCA already guarantees GPU/device
+    // scope, so only add a system release when the caller required system scope
+    // but gave a weaker (numerically larger) scope: given > required. Do NOT
+    // "align" this with the HIP backends' given < required -- that would make
+    // this fence dead code here.
     if ((required == cuda::thread_scope_system) && (given > required)) {
       doca_gpu_dev_verbs_fence_release<DOCA_GPUNETIO_VERBS_SYNC_SCOPE_SYS>();
     }
@@ -181,8 +186,13 @@ NCCL_DEVICE_INLINE static void putValueImplMode(
       sig_laddr.key = loadConst(&gdaki->sink_buffer_lkey);
     }
 
-    // cuda::thread_scope_system has the lowest value
-    // DOCA guarantees SCOPE_GPU. Only add another release if SCOPE_SYSTEM is required.
+    // gdaki is CUDA-only (NCCL_GIN_GDAKI_ENABLE=0 on HIP). In libcu++,
+    // cuda::thread_scope_system is the LOWEST value -- the opposite of HIP's
+    // hip_compat.h, where it is the max. DOCA already guarantees GPU/device
+    // scope, so only add a system release when the caller required system scope
+    // but gave a weaker (numerically larger) scope: given > required. Do NOT
+    // "align" this with the HIP backends' given < required -- that would make
+    // this fence dead code here.
     if ((required == cuda::thread_scope_system) && (given > required)) {
       doca_gpu_dev_verbs_fence_release<DOCA_GPUNETIO_VERBS_SYNC_SCOPE_SYS>();
     }

projects/rccl/src/include/nccl_device/gin/gin_device_common.h

@@ -56,6 +56,14 @@
 #endif
 #endif
 
+// Test seam for the system-scope fence issued by the HIP GIN Put/PutValue
+// templates. Expands to the real fence in production (zero cost); unit tests may
+// override it (before including a backend header) to observe that the fence
+// branch actually executed.
+#ifndef NCCL_GIN_THREADFENCE_SYSTEM
+#define NCCL_GIN_THREADFENCE_SYSTEM() __threadfence_system()
+#endif
+
 enum ncclGinOptFlags {
   ncclGinOptFlagsDefault = 0,
   ncclGinOptFlagsMaySkipCreditCheck = (1 << 0),

projects/rccl/src/include/nccl_device/gin/rocshmem_gda/gin_rocshmem_gda.h

@@ -29,8 +29,10 @@ struct ncclGinApi_Put<NCCL_NET_DEVICE_GIN_ROCSHMEM_GDA> {
       rocshmem::QueuePair* qp = loadConst(loadConst(&rsCtx->qps) + peer);
       rocshmem::ActiveWFInfo wf_info(peer, rocshmem::ThreadScope::thread);
 
-      if ((required == cuda::thread_scope_system) && (given > required)) {
-        __threadfence_system();
+      // HIP thread_scope (hip_compat.h): system is the MAX value, so a caller that
+      // only guaranteed a weaker scope has given < required -> add a system fence.
+      if ((required == cuda::thread_scope_system) && (given < required)) {
+        NCCL_GIN_THREADFENCE_SYSTEM();
       }
 
       // Skip zero-length RDMA writes (0-byte put_nbi can stall quiet()/flush() -> deadlock); signal still delivered, matching native rocSHMEM/PROXY.
@@ -85,8 +87,10 @@ struct ncclGinApi_PutValue<NCCL_NET_DEVICE_GIN_ROCSHMEM_GDA> {
       uintptr_t dstAddr = loadConst(loadConst(&dstMh->remote_vas) + peer) + dstOff;
       uint32_t dstRkey = loadConst(loadConst(&dstMh->rkeys) + peer);
 
-      if ((required == cuda::thread_scope_system) && (given > required)) {
-        __threadfence_system();
+      // HIP thread_scope (hip_compat.h): system is the MAX value, so a caller that
+      // only guaranteed a weaker scope has given < required -> add a system fence.
+      if ((required == cuda::thread_scope_system) && (given < required)) {
+        NCCL_GIN_THREADFENCE_SYSTEM();
       }
 
       // lkey=0: put_nbi copies srcVal inline into the WQE

projects/rccl/test/device/GinAnvilSdmaTemplate_test.cpp

@@ -15,6 +15,11 @@
 #include "nccl_device/gin/anvil_sdma/gin_anvil_sdma_device_host_common.h"
 
 #if NCCL_GIN_ANVIL_SDMA_ENABLE
+// Count invocations of the Put/PutValue system-scope fence seam (gin_device_common.h).
+// Override must precede gin_anvil_sdma.h so the templates expand our counter.
+__device__ unsigned long long g_sdmaStubThreadfenceCount = 0;
+#undef NCCL_GIN_THREADFENCE_SYSTEM
+#define NCCL_GIN_THREADFENCE_SYSTEM() atomicAdd(&g_sdmaStubThreadfenceCount, 1ULL)
 #include "nccl_device/gin/anvil_sdma/gin_anvil_sdma.h"
 #endif
 
@@ -68,6 +73,17 @@ static void uploadHarness(DeviceBuffer<TemplateHarness>* d_h, TemplateHarness* h
   d_h->upload(*host);
 }
 
+static void resetThreadfenceCount() {
+  unsigned long long z = 0;
+  HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(g_sdmaStubThreadfenceCount), &z, sizeof(z)));
+}
+
+static unsigned long long readThreadfenceCount() {
+  unsigned long long c = 0;
+  HIP_EXPECT(hipMemcpyFromSymbol(&c, HIP_SYMBOL(g_sdmaStubThreadfenceCount), sizeof(c)));
+  return c;
+}
+
 // H1: non-leader thread returns immediately.
 __global__ void kernelPutLeaderOnly(TemplateHarness* h, int* executed) {
   ncclGinCtx ginCtx{};
@@ -98,7 +114,7 @@ TEST_F(GinAnvilSdmaTemplateTest, Put_NonLeaderThreadNoOp) {
   EXPECT_EQ(d_executed.download(), 1);
 }
 
-// H2: thread-scope fence when given > required.
+// H2: system fence when required==system && given<required (HIP scope ordering).
 __global__ void kernelPutScopeFence(TemplateHarness* h) {
   if (threadIdx.x != 0) return;
   ncclGinCtx ginCtx{};
@@ -121,8 +137,10 @@ TEST_F(GinAnvilSdmaTemplateTest, Put_ThreadScopeFence) {
   DeviceBuffer<TemplateHarness> d_h(1);
   TemplateHarness host{};
   uploadHarness(&d_h, &host, &d_src, &d_dst, &d_entry, &d_q, &d_row, 128);
+  resetThreadfenceCount();
   kernelPutScopeFence<<<1, 1>>>(d_h.ptr);
   syncAndCheck();
+  EXPECT_EQ(readThreadfenceCount(), 1ULL);
 }
 
 // H3: SDMA path (threshold 0) with stub put + markSdmaDirty.

projects/rccl/test/device/GinRocshmemGdaTemplate_test.cpp

@@ -24,6 +24,11 @@
 #include "nccl_device/gin/rocshmem_gda/gin_rocshmem_device_host_common_gda.h"
 
 #if NCCL_GIN_ROCSHMEM_GDA_ENABLE
+// Count invocations of the Put/PutValue system-scope fence seam (gin_device_common.h).
+// Override must precede gin_rocshmem_gda.h so the templates expand our counter.
+__device__ unsigned long long g_gdaStubThreadfenceCount = 0;
+#undef NCCL_GIN_THREADFENCE_SYSTEM
+#define NCCL_GIN_THREADFENCE_SYSTEM() atomicAdd(&g_gdaStubThreadfenceCount, 1ULL)
 #include "nccl_device/gin/rocshmem_gda/gin_rocshmem_gda.h"
 #endif
 
@@ -172,6 +177,17 @@ static unsigned long long readPutNbiCount() {
   return n;
 }
 
+static void resetThreadfenceCount() {
+  unsigned long long z = 0;
+  HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(g_gdaStubThreadfenceCount), &z, sizeof(z)));
+}
+
+static unsigned long long readThreadfenceCount() {
+  unsigned long long c = 0;
+  HIP_EXPECT(hipMemcpyFromSymbol(&c, HIP_SYMBOL(g_gdaStubThreadfenceCount), sizeof(c)));
+  return c;
+}
+
 // G1: Put with data (no signal) copies src -> peer's remote buffer.
 __global__ void kernelPutData(GdaHarness* h, size_t bytes) {
   ncclGinCtx ginCtx{};
@@ -336,12 +352,7 @@ TEST_F(GinRocshmemGdaTemplateTest, Put_SignalAndCounter) {
   EXPECT_EQ(ctr[1], 1ULL);
 }
 
-// G7: weaker-given-scope path (required=system, given=block) still completes the
-// put. NOTE: the template's fence guard is `(required==system && given>required)`;
-// since `system` is the maximum cuda::thread_scope, that branch is never taken
-// (here or anywhere). This guard is a pre-existing convention shared by all GIN
-// backends (anvil_sdma, gdaki, rocshmem_gda) and is tracked for a separate,
-// coordinated fix -- so this case only asserts the data lands, not that a fence ran.
+// G7: required=system, given=block -> HIP guard fires (given < required) and put completes.
 __global__ void kernelPutScopeFence(GdaHarness* h) {
   ncclGinCtx ginCtx{};
   ginCtx.handle = &h->ctx;
@@ -354,16 +365,35 @@ __global__ void kernelPutScopeFence(GdaHarness* h) {
       false, nullptr, cuda::thread_scope_system, cuda::thread_scope_block);
 }
 
-TEST_F(GinRocshmemGdaTemplateTest, Put_WeakerGivenScopeStillPuts) {
+TEST_F(GinRocshmemGdaTemplateTest, Put_WeakerGivenScopeFencesAndPuts) {
   constexpr int kN = 8;
   std::vector<uint8_t> pat(kN);
   for (int i = 0; i < kN; ++i) pat[static_cast<size_t>(i)] = static_cast<uint8_t>(0x11 * (i + 1));
   GdaEnv env(kN);
   env.src.copyFrom(pat);
   env.dst.zero();
   env.build();
+  resetThreadfenceCount();
   kernelPutScopeFence<<<1, 1>>>(env.dHarness.ptr);
   syncAndCheck();
+  EXPECT_EQ(readThreadfenceCount(), 1ULL);
+  auto got = env.dst.copyTo();
+  for (int i = 0; i < kN; ++i) EXPECT_EQ(got[static_cast<size_t>(i)], pat[static_cast<size_t>(i)]);
+}
+
+// G7b: required=given=system -> guard does not fire (given < required is false) and put completes.
+TEST_F(GinRocshmemGdaTemplateTest, Put_EqualScopeTakesNoFence) {
+  constexpr int kN = 8;
+  std::vector<uint8_t> pat(kN);
+  for (int i = 0; i < kN; ++i) pat[static_cast<size_t>(i)] = static_cast<uint8_t>(0x22 + i);
+  GdaEnv env(kN);
+  env.src.copyFrom(pat);
+  env.dst.zero();
+  env.build();
+  resetThreadfenceCount();
+  kernelPutData<<<1, 1>>>(env.dHarness.ptr, kN);
+  syncAndCheck();
+  EXPECT_EQ(readThreadfenceCount(), 0ULL);
   auto got = env.dst.copyTo();
   for (int i = 0; i < kN; ++i) EXPECT_EQ(got[static_cast<size_t>(i)], pat[static_cast<size_t>(i)]);
 }