diff --git a/projects/rccl/src/algorithms/gin/sdma/gin_all_reduce_sdma.cu b/projects/rccl/src/algorithms/gin/sdma/gin_all_reduce_sdma.cu index 988a566131ba..f62f7aa978b6 100755 --- a/projects/rccl/src/algorithms/gin/sdma/gin_all_reduce_sdma.cu +++ b/projects/rccl/src/algorithms/gin/sdma/gin_all_reduce_sdma.cu @@ -67,7 +67,8 @@ static ncclResult_t ncclGinAllReduceInitOnce(ncclComm* comm) { reqs.barrierCount = kGinAllReduceLsaCtas; reqs.ginSignalCount = kGinAllReduceLsaCtas; reqs.ginConnectionType = NCCL_GIN_CONNECTION_FULL; - NCCLCHECK(ncclDevrCommCreateInternal(comm, &reqs, &state->devComm, /*isInternal=*/true)); + NCCLCHECK(ncclDevrCommCreateInternal(comm, &reqs, &state->devComm, /*isInternal=*/true, + /*deviceCodeVersion=*/NCCL_VERSION_CODE)); state->initialized = true; // Two words (arrived, sense) for ginIntraGpuCtaBarrier. ncclCudaCalloc is capture-safe @@ -231,7 +232,9 @@ static bool ginAllReduceBaseEligible(ncclComm* comm, const void* sendbuff, void* if (comm->nNodes != 1) return false; if (comm->nRanks > kGinAllReduceMaxRanks) return false; if (comm->sharedRes == nullptr) return false; - if (comm->sharedRes->ginState.ginType != (ncclGinType_t)NCCL_NET_DEVICE_GIN_ANVIL_SDMA) return false; + ncclGinType_t ginType = NCCL_GIN_TYPE_NONE; + if (ncclGetGinType(comm, &ginType) != ncclSuccess) return false; + if (ginType != (ncclGinType_t)NCCL_NET_DEVICE_GIN_ANVIL_SDMA) return false; // After cheaper ginType: ncclTeamLsa() can initialize the device runtime. if (ncclTeamLsa(comm).nRanks != comm->nRanks) return false; return true; diff --git a/projects/rccl/src/algorithms/gin/sdma/gin_alltoall_sdma.cu b/projects/rccl/src/algorithms/gin/sdma/gin_alltoall_sdma.cu index 35a4fd8dc720..2a160baccd01 100644 --- a/projects/rccl/src/algorithms/gin/sdma/gin_alltoall_sdma.cu +++ b/projects/rccl/src/algorithms/gin/sdma/gin_alltoall_sdma.cu @@ -158,7 +158,8 @@ ncclResult_t ncclGinA2AInitOnce(ncclComm* comm) { reqs.ginSignalCount = kGinA2AMaxCtas; reqs.barrierCount = kGinA2ASdmaCtas; reqs.ginConnectionType = NCCL_GIN_CONNECTION_FULL; - NCCLCHECK(ncclDevrCommCreateInternal(comm, &reqs, &state->devComm, /*isInternal=*/true)); + NCCLCHECK(ncclDevrCommCreateInternal(comm, &reqs, &state->devComm, /*isInternal=*/true, + /*deviceCodeVersion=*/NCCL_VERSION_CODE)); state->initialized = true; } return ncclSuccess; @@ -181,7 +182,9 @@ bool ncclAllToAllGinSdmaEligible(ncclComm* comm, const void* sendbuff, void* rec if (comm->globalGinSupport != NCCL_GIN_CONNECTION_FULL) return false; // This path runs on the comm's shared GIN backend, so it has to be SDMA. - if (comm->sharedRes->ginState.ginType != (ncclGinType_t)NCCL_NET_DEVICE_GIN_ANVIL_SDMA) return false; + ncclGinType_t ginType = NCCL_GIN_TYPE_NONE; + if (ncclGetGinType(comm, &ginType) != ncclSuccess) return false; + if (ginType != (ncclGinType_t)NCCL_NET_DEVICE_GIN_ANVIL_SDMA) return false; // Every rank must be reachable over LSA. if (ncclTeamLsa(comm).nRanks != comm->nRanks) return false; 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 c279ae72d142..c575c0d0bd29 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 @@ -469,4 +469,13 @@ struct ncclGinApi_Wait { } }; +// Only Flush drains the SDMA engine, so the barrier has to keep its pre-signal +// fenceFlush. +template <> +struct ncclGinApi_SupportsStrongSignal { + NCCL_DEVICE_INLINE static bool call(ncclGinCtx) { + return false; + } +}; + #endif diff --git a/projects/rccl/src/include/nccl_device/gin/rocshmem_gda/gin_rocshmem_gda.h b/projects/rccl/src/include/nccl_device/gin/rocshmem_gda/gin_rocshmem_gda.h index 74bc8d19e688..168db592cef8 100644 --- a/projects/rccl/src/include/nccl_device/gin/rocshmem_gda/gin_rocshmem_gda.h +++ b/projects/rccl/src/include/nccl_device/gin/rocshmem_gda/gin_rocshmem_gda.h @@ -186,4 +186,12 @@ struct ncclGinApi_Wait { } }; +// Mirrors ginRocshmemGdaGetGinProperties()'s supportsStrongSignals on the host side. +template <> +struct ncclGinApi_SupportsStrongSignal { + NCCL_DEVICE_INLINE static bool call(ncclGinCtx) { + return true; + } +}; + #endif /* _NCCL_DEVICE_GIN_ROCSHMEM_GDA_H_ */ diff --git a/projects/rccl/test/GinAlltoAllEligibilityTests.cpp b/projects/rccl/test/GinAlltoAllEligibilityTests.cpp index d5586e7bcc33..59e215cc3b16 100644 --- a/projects/rccl/test/GinAlltoAllEligibilityTests.cpp +++ b/projects/rccl/test/GinAlltoAllEligibilityTests.cpp @@ -44,8 +44,9 @@ struct GinAlltoAllMockComm comm.devrState.lsaSize = comm.nRanks; comm.devrState.lsaSelf = comm.rank; - sharedRes.ginState.ginType = (ncclGinType_t)NCCL_NET_DEVICE_GIN_ANVIL_SDMA; - comm.sharedRes = &sharedRes; + sharedRes.ginState.numActiveBackends = 1; + sharedRes.ginState.backends[0].ginType = (ncclGinType_t)NCCL_NET_DEVICE_GIN_ANVIL_SDMA; + comm.sharedRes = &sharedRes; } ncclComm* get() { return &comm; } @@ -146,13 +147,23 @@ TEST_F(GinAlltoAllEligibilityTest, LsaTeamSmallerThanComm) // other backend must not be claimed. TEST_F(GinAlltoAllEligibilityTest, WrongGinBackend) { - mockComm_.sharedRes.ginState.ginType = (ncclGinType_t)NCCL_NET_DEVICE_GIN_ROCSHMEM_GDA; + mockComm_.sharedRes.ginState.backends[0].ginType = (ncclGinType_t)NCCL_NET_DEVICE_GIN_ROCSHMEM_GDA; EXPECT_FALSE(eligible(kCount)); - mockComm_.sharedRes.ginState.ginType = (ncclGinType_t)NCCL_NET_DEVICE_GIN_PROXY; + mockComm_.sharedRes.ginState.backends[0].ginType = (ncclGinType_t)NCCL_NET_DEVICE_GIN_PROXY; EXPECT_FALSE(eligible(kCount)); } +// WrongGinBackend above cannot pin which backends[] slot the gate reads: the empty +// window registry rejects every case anyway, so writing the wrong index would still +// leave it passing. Assert the accessor directly against the fixture's default state. +TEST_F(GinAlltoAllEligibilityTest, MockBackendSlotMatchesGinTypeAccessor) +{ + ncclGinType_t ginType = NCCL_GIN_TYPE_NONE; + ASSERT_EQ(ncclGetGinType(mockComm_.get(), &ginType), ncclSuccess); + EXPECT_EQ(ginType, (ncclGinType_t)NCCL_NET_DEVICE_GIN_ANVIL_SDMA); +} + // The kernel puts straight into symmetric windows, so unregistered buffers are // rejected. This is the mock's default state: every gate above passes, and the // empty window registry is the only reason the verdict is false. diff --git a/projects/rccl/tools/ci/lib/gin-tests.json b/projects/rccl/tools/ci/lib/gin-tests.json index 66e91b841356..71f623a8fe67 100644 --- a/projects/rccl/tools/ci/lib/gin-tests.json +++ b/projects/rccl/tools/ci/lib/gin-tests.json @@ -44,7 +44,7 @@ "name": "rccl-gin-rocshmem-sdma-bcast-d3", "kind": "rccl-tests", "bin": "broadcast_perf", - "env": ["NCCL_CUMEM_ENABLE=1", "RCCL_ENABLE_INTRANET=1", "NCCL_DMABUF_ENABLE=1", "HSA_NO_SCRATCH_RECLAIM=1", "NCCL_MSCCL_ENABLE=0", "NCCL_GIN_ENABLE=1", "NCCL_GIN_TYPE=6"], + "env": ["NCCL_CUMEM_ENABLE=1", "RCCL_ENABLE_INTRANET=1", "NCCL_DMABUF_ENABLE=1", "HSA_NO_SCRATCH_RECLAIM=1", "NCCL_MSCCL_ENABLE=0", "NCCL_GIN_ENABLE=1", "NCCL_GIN_TYPE=7"], "args": "-b 128 -e {E} -f 2 -g 1 -R 2 -D 3 -A 1" }, { @@ -72,7 +72,7 @@ "name": "rccl-gin-rocshmem-sdma-internal", "kind": "rccl-tests", "bin": "alltoall_perf", - "env": ["NCCL_CUMEM_ENABLE=1", "NCCL_DMABUF_ENABLE=1", "HSA_NO_SCRATCH_RECLAIM=1", "NCCL_GIN_ENABLE=1", "NCCL_GIN_TYPE=6"], + "env": ["NCCL_CUMEM_ENABLE=1", "NCCL_DMABUF_ENABLE=1", "HSA_NO_SCRATCH_RECLAIM=1", "NCCL_GIN_ENABLE=1", "NCCL_GIN_TYPE=7"], "args": "-b 128 -e {E} -f 2 -g 1 -R 2" } ] diff --git a/projects/rccl/tools/ci/run-gin-ci.sh b/projects/rccl/tools/ci/run-gin-ci.sh index 00a610c3cd07..35dc8e294f52 100755 --- a/projects/rccl/tools/ci/run-gin-ci.sh +++ b/projects/rccl/tools/ci/run-gin-ci.sh @@ -22,7 +22,7 @@ # GIN_PYTEST_TIMEOUT Wall-clock cap for pytest matrix entries (default: 1800s) # GIN_PYTEST_HW_CASES mpirun cases under -k GinSdma (default: 9). Offline # parser/tier guards do not launch and are not counted. -# RCCL_TESTS_BCAST_GIN_TYPE NCCL_GIN_TYPE for Broadcast pytest (default: 6) +# RCCL_TESTS_BCAST_GIN_TYPE NCCL_GIN_TYPE for Broadcast pytest (default: 7) # RCCL_TESTS_BCAST_TIMEOUT_S / RCCL_TESTS_BCAST_CONN_RETRIES # Inner pytest launch budget. Unset → derived so # HW_CASES * retries * TIMEOUT_S is strictly under @@ -245,7 +245,7 @@ run_test() { RCCL_TESTS_GIN_SDMA_BCAST=1 \ RCCL_TESTS_BCAST_EXE="${bcast_exe}" \ RCCL_TESTS_BCAST_NP="${NP}" \ - RCCL_TESTS_BCAST_GIN_TYPE="${RCCL_TESTS_BCAST_GIN_TYPE:-6}" \ + RCCL_TESTS_BCAST_GIN_TYPE="${RCCL_TESTS_BCAST_GIN_TYPE:-7}" \ RCCL_TESTS_BCAST_TIMEOUT_S="${RCCL_TESTS_BCAST_TIMEOUT_S}" \ RCCL_TESTS_BCAST_CONN_RETRIES="${RCCL_TESTS_BCAST_CONN_RETRIES}" \ RCCL_TESTS_BCAST_XENV="${bcast_xenv}" \ diff --git a/projects/rccl/tools/scripts/test_runner/configs/gin_rocshmem_backends_mellanox.json b/projects/rccl/tools/scripts/test_runner/configs/gin_rocshmem_backends_mellanox.json index 0a8ea9508f28..123cf705735a 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/gin_rocshmem_backends_mellanox.json +++ b/projects/rccl/tools/scripts/test_runner/configs/gin_rocshmem_backends_mellanox.json @@ -55,7 +55,7 @@ "gin_perf_gda": { "extends": "gin_perf_base", "env_variables": { - "NCCL_GIN_TYPE": "5" + "NCCL_GIN_TYPE": "6" }, "tests": [ { "name": "GIN_GDA_D3", "description": "GDA alltoall, GinAlltoAllKernel, pure-GIN (-R 2 -D 3)", "command_args": "-R 2 -D 3" }, @@ -68,7 +68,7 @@ "num_ranks": 8, "num_nodes": 1, "env_variables": { - "NCCL_GIN_TYPE": "5" + "NCCL_GIN_TYPE": "6" }, "tests": [ { "name": "GIN_GDA_1N_D3", "description": "GDA alltoall single-node, GinAlltoAllKernel, pure-GIN (-R 2 -D 3)", "command_args": "-R 2 -D 3" }, @@ -81,7 +81,7 @@ "num_ranks": 8, "num_nodes": 1, "env_variables": { - "NCCL_GIN_TYPE": "6" + "NCCL_GIN_TYPE": "7" }, "tests": [ { "name": "GIN_SDMA_D3", "description": "SDMA alltoall, GinAlltoAllKernel, pure-GIN (-R 2 -D 3)", "command_args": "-R 2 -D 3" }, @@ -95,7 +95,7 @@ "num_nodes": 1, "command_args": "-b 1 -e 4G -f 2 -g 1", "env_variables": { - "NCCL_GIN_TYPE": "6" + "NCCL_GIN_TYPE": "7" }, "tests": [ { "name": "GIN_SDMA_INTERNAL", "description": "GIN-SDMA alltoall via ncclAlltoAll (-R 2)", "command_args": "-R 2" } @@ -191,7 +191,7 @@ "gin_devmpi_gda": { "extends": "gin_devmpi_common", "env_variables": { - "NCCL_GIN_TYPE": "5" + "NCCL_GIN_TYPE": "6" } }, "gin_devmpi_sdma": { @@ -200,7 +200,7 @@ "num_nodes": 1, "num_gpus": 2, "env_variables": { - "NCCL_GIN_TYPE": "6" + "NCCL_GIN_TYPE": "7" }, "tests": [ { "name": "BarrierSession_LsaOnly", "num_gpus": 4 },