refactor(rccl/gin): fold Anvil Flush drain and BarrierFence test helpers

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

@@ -282,15 +282,15 @@ NCCL_DEVICE_INLINE void quietOwnedSdmaDirtyBits(ncclGinAnvilSdmaGPUContext* rsCt
     __scoped_atomic_fetch_and(sdmaDirty, ~owned, __ATOMIC_RELAXED, __MEMORY_SCOPE_DEVICE);
     return;
   }
+  // Bit p * numCh + ch is also that queue's index in handles[], and owned has
+  // already dropped bits past kSdmaDirtyBitWidth.
+  uint64_t rem = dirty;
 #pragma unroll 1
-  for (int p = threadRank; p < nRanks; p += stride) {
-    for (int ch = 0; ch < numCh; ++ch) {
-      const int bitIdx = p * numCh + ch;
-      if (bitIdx >= kSdmaDirtyBitWidth) break;
-      if ((dirty & (1ULL << bitIdx)) == 0) continue;
-      auto* h = loadConst(handles + bitIdx);
-      if (h != nullptr) ::sdma_anvil::quiet(*h);
-    }
+  while (rem != 0) {
+    const int bitIdx = __builtin_ctzll(rem);
+    rem &= rem - 1;
+    auto* h = loadConst(handles + bitIdx);
+    if (h != nullptr) ::sdma_anvil::quiet(*h);
   }
   __scoped_atomic_fetch_and(sdmaDirty, ~dirty, __ATOMIC_RELAXED, __MEMORY_SCOPE_DEVICE);
 }

projects/rccl/test/transport/GinDeviceMPITests.cpp

@@ -2040,29 +2040,31 @@ TEST_F(GinMPIDeviceTests, WaitCounterAndSignal_MultiContext) {
   runWaitCounterAndSignal(n);
 }
 
+// Fill and check byte for BarrierFence visibility: 0x20 + rank + (i & 0x3f).
+__host__ __device__ inline uint8_t barrierFencePattern(int rank, size_t i) {
+  return static_cast<uint8_t>(0x20 + rank + (i & 0x3f));
+}
+
+template <typename Coop, typename GinOrAllContexts>
+__device__ void syncFenceVisibilityBarrierWith(Coop coop, GinOrAllContexts ginOrAllContexts,
+                                               BarrierFenceOperation operation, bool defaultFence) {
+  ncclGinBarrierSession<Coop> bar{coop, ginOrAllContexts, ncclTeamTagWorld{}, /*barrierIndex=*/0};
+  if (defaultFence) {
+    bar.sync(coop, cuda::memory_order_acq_rel);
+  } else {
+    ncclGinFenceLevel fence =
+      operation == BarrierFenceOperation::Get ? ncclGinFenceLevel::Get : ncclGinFenceLevel::Put;
+    bar.sync(coop, cuda::memory_order_acq_rel, fence);
+  }
+}
+
 template <typename GinOrAllContexts>
 __device__ void syncFenceVisibilityBarrier(GinOrAllContexts ginOrAllContexts, BarrierFenceOperation operation,
                                            bool defaultFence, bool coopAny) {
   if (coopAny) {
-    ncclCoopAny coop{ncclCoopCta()};
-    ncclGinBarrierSession<ncclCoopAny> bar{coop, ginOrAllContexts, ncclTeamTagWorld{}, /*barrierIndex=*/0};
-    if (defaultFence) {
-      bar.sync(coop, cuda::memory_order_acq_rel);
-    } else {
-      ncclGinFenceLevel fence =
-        operation == BarrierFenceOperation::Get ? ncclGinFenceLevel::Get : ncclGinFenceLevel::Put;
-      bar.sync(coop, cuda::memory_order_acq_rel, fence);
-    }
+    syncFenceVisibilityBarrierWith(ncclCoopAny{ncclCoopCta()}, ginOrAllContexts, operation, defaultFence);
   } else {
-    ncclGinBarrierSession<ncclCoopCta> bar{
-      ncclCoopCta(), ginOrAllContexts, ncclTeamTagWorld{}, /*barrierIndex=*/0};
-    if (defaultFence) {
-      bar.sync(ncclCoopCta(), cuda::memory_order_acq_rel);
-    } else {
-      ncclGinFenceLevel fence =
-        operation == BarrierFenceOperation::Get ? ncclGinFenceLevel::Get : ncclGinFenceLevel::Put;
-      bar.sync(ncclCoopCta(), cuda::memory_order_acq_rel, fence);
-    }
+    syncFenceVisibilityBarrierWith(ncclCoopCta(), ginOrAllContexts, operation, defaultFence);
   }
 }
 
@@ -2093,8 +2095,7 @@ __global__ void barrierFenceVisibilityKernel(
 
   int sourceRank = operation == BarrierFenceOperation::Put ? (rank + devComm.nRanks - 1) % devComm.nRanks : peer;
   for (size_t i = threadIdx.x; i < bytes; i += blockDim.x) {
-    uint8_t expected = static_cast<uint8_t>(0x20 + sourceRank + (i & 0x3f));
-    if (dst[i] != expected) atomicCAS(error, 0, static_cast<int>(i + 1));
+    if (dst[i] != barrierFencePattern(sourceRank, i)) atomicCAS(error, 0, static_cast<int>(i + 1));
   }
 }
 
@@ -2117,8 +2118,7 @@ __global__ void barrierFenceIpcSignalIncKernel(
   }
   gin.waitSignal(ncclCoopCta(), kSigIdx, /*least=*/1);
   for (size_t i = threadIdx.x; i < bytes; i += blockDim.x) {
-    uint8_t expected = static_cast<uint8_t>(0x20 + /*sourceRank=*/0 + (i & 0x3f));
-    if (dst[i] != expected) atomicCAS(error, 0, static_cast<int>(i + 1));
+    if (dst[i] != barrierFencePattern(/*sourceRank=*/0, i)) atomicCAS(error, 0, static_cast<int>(i + 1));
   }
 }
 
@@ -2182,6 +2182,9 @@ GinSymmetricSetupStatus prepareGinSymmetricBuffers(GinSymmetricBuffers* b, ncclC
   return st;
 }
 
+// ASSERT_MPI_EQ runs in the caller, so a peer-skip GTEST_SKIP returns from the
+// test body. Checking IsSkipped() after a helper that asserted internally would
+// miss ranks that already have an EXPECT failure (IsSkipped() is then false).
 #define ASSERT_GIN_SYMMETRIC_SETUP(st)     \
   do {                                     \
     SCOPED_TRACE((st).step);               \
@@ -2214,8 +2217,9 @@ void GinMPIDeviceTests::runBarrierFenceVisibility(
   // exercises a real SDMA copy rather than the small-transfer IPC fallback.
   constexpr size_t kBytes = 4096;
   std::vector<uint8_t> hostSrc(kBytes);
-  for (size_t i = 0; i < kBytes; ++i)
-    hostSrc[i] = static_cast<uint8_t>(0x20 + rank + (i & 0x3f));
+  for (size_t i = 0; i < kBytes; ++i) {
+    hostSrc[i] = barrierFencePattern(rank, i);
+  }
 
   ncclDevCommRequirements reqs = defaultGinReqs();
   reqs.worldGinBarrierCount = 1;
@@ -2298,8 +2302,9 @@ TEST_F(GinMPIDeviceTests, BarrierFence_IpcSignalIncMakesPutVisible_SingleNode) {
 
   constexpr size_t kBytes = 64;
   std::vector<uint8_t> hostSrc(kBytes);
-  for (size_t i = 0; i < kBytes; ++i)
-    hostSrc[i] = static_cast<uint8_t>(0x20 + /*rank0 pattern*/0 + (i & 0x3f));
+  for (size_t i = 0; i < kBytes; ++i) {
+    hostSrc[i] = barrierFencePattern(/*rank0 pattern*/ 0, i);
+  }
 
   ncclDevCommRequirements reqs = defaultGinReqs();
   reqs.worldGinBarrierCount = 1;

projects/rccl/tools/ci/run-gin-ci.sh

@@ -203,8 +203,7 @@ run_test() {
   case "${kind}" in
     rocshmem)     bin_path="${ROCSHMEM_TESTS_BIN_DIR}/${bin}" ;;
     rccl-tests)   bin_path="${RCCL_TESTS_BIN_DIR}/${bin}" ;;
-    fixtures)     bin_path="${RCCL_FIXTURES_BIN_DIR}/${bin}" ;;
-    mpi-fixtures) bin_path="${RCCL_FIXTURES_BIN_DIR}/${bin}" ;;
+    fixtures|mpi-fixtures) bin_path="${RCCL_FIXTURES_BIN_DIR}/${bin}" ;;
     pytest)       bench_timeout="${GIN_PYTEST_TIMEOUT}" ;;
     *) echo "  SKIP ${name}: unknown kind '${kind}'"; FAILED_RUNS+=("${name} (unknown kind)"); return ;;
   esac
@@ -260,16 +259,15 @@ run_test() {
     timeout --kill-after="${BENCH_KILL_AFTER}" "${bench_timeout}" \
       env LD_LIBRARY_PATH="${LD_LIBRARY_PATH}" \
         "${bin_path}" ${args}
-  elif [[ "${kind}" == "mpi-fixtures" ]]; then
+  else
+    local ld_flag="LD_LIBRARY_PATH"
     # rccl-UnitTestsMPI links librccl internals, which only the build tree's
     # ENABLE_MPI_TESTS librccl exports; the installed copy hides them.
+    if [[ "${kind}" == "mpi-fixtures" ]]; then
+      ld_flag="LD_LIBRARY_PATH=${RCCL_BUILD_LIB_DIR}:${LD_LIBRARY_PATH}"
+    fi
     timeout --kill-after="${BENCH_KILL_AFTER}" "${bench_timeout}" \
-      mpirun -np "${nranks}" ${MCA} ${env_flags} \
-        -x LD_LIBRARY_PATH="${RCCL_BUILD_LIB_DIR}:${LD_LIBRARY_PATH}" \
-        "${bin_path}" ${args}
-  else
-    timeout --kill-after="${BENCH_KILL_AFTER}" "${bench_timeout}" \
-      mpirun -np "${nranks}" ${MCA} ${env_flags} -x LD_LIBRARY_PATH \
+      mpirun -np "${nranks}" ${MCA} ${env_flags} -x "${ld_flag}" \
         "${bin_path}" ${args}
   fi
   local rc=$?