diff --git a/projects/rccl/src/include/nccl_device/gin/anvil_sdma/gin_anvil_sdma.h b/projects/rccl/src/include/nccl_device/gin/anvil_sdma/gin_anvil_sdma.h index 215286f52562..3fb01b34bccb 100644 --- a/projects/rccl/src/include/nccl_device/gin/anvil_sdma/gin_anvil_sdma.h +++ b/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); } diff --git a/projects/rccl/test/transport/GinDeviceMPITests.cpp b/projects/rccl/test/transport/GinDeviceMPITests.cpp index 753aff05fb5f..bd4544c17673 100644 --- a/projects/rccl/test/transport/GinDeviceMPITests.cpp +++ b/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(0x20 + rank + (i & 0x3f)); +} + +template +__device__ void syncFenceVisibilityBarrierWith(Coop coop, GinOrAllContexts ginOrAllContexts, + BarrierFenceOperation operation, bool defaultFence) { + ncclGinBarrierSession 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 __device__ void syncFenceVisibilityBarrier(GinOrAllContexts ginOrAllContexts, BarrierFenceOperation operation, bool defaultFence, bool coopAny) { if (coopAny) { - ncclCoopAny coop{ncclCoopCta()}; - ncclGinBarrierSession 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 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(0x20 + sourceRank + (i & 0x3f)); - if (dst[i] != expected) atomicCAS(error, 0, static_cast(i + 1)); + if (dst[i] != barrierFencePattern(sourceRank, i)) atomicCAS(error, 0, static_cast(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(0x20 + /*sourceRank=*/0 + (i & 0x3f)); - if (dst[i] != expected) atomicCAS(error, 0, static_cast(i + 1)); + if (dst[i] != barrierFencePattern(/*sourceRank=*/0, i)) atomicCAS(error, 0, static_cast(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 hostSrc(kBytes); - for (size_t i = 0; i < kBytes; ++i) - hostSrc[i] = static_cast(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 hostSrc(kBytes); - for (size_t i = 0; i < kBytes; ++i) - hostSrc[i] = static_cast(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; diff --git a/projects/rccl/tools/ci/run-gin-ci.sh b/projects/rccl/tools/ci/run-gin-ci.sh index 55b68062f31b..bb77556548b5 100755 --- a/projects/rccl/tools/ci/run-gin-ci.sh +++ b/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=$?