Skip to content

fix(rccl): Fix SDMA Collectives build failure after Sync #11765

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
Kapil-Shyam-Pawar merged 5 commits into from
Sep 18, 2026
Merged

fix(rccl): Fix SDMA Collectives build failure after Sync #11765

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
b212014
Fix build failure
Kapil-Shyam-Pawar Sep 17, 2026
25e5065
Update gin type for broadcast
Kapil-Shyam-Pawar Sep 17, 2026
4a8d2cd
Update test
Kapil-Shyam-Pawar Sep 17, 2026
77063ae
Fix test build failure and update config
Kapil-Shyam-Pawar Sep 17, 2026
fc041d8
Update config
Kapil-Shyam-Pawar Sep 18, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
7 changes: 5 additions & 2 deletions projects/rccl/src/algorithms/gin/sdma/gin_all_reduce_sdma.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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));
Comment thread
Kapil-Shyam-Pawar marked this conversation as resolved.
Resolved
state->initialized = true;
Comment thread
Kapil-Shyam-Pawar marked this conversation as resolved.
Resolved

// Two words (arrived, sense) for ginIntraGpuCtaBarrier. ncclCudaCalloc is capture-safe
Expand Down Expand Up @@ -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;
Comment thread
Kapil-Shyam-Pawar marked this conversation as resolved.
Resolved
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;
Expand Down
7 changes: 5 additions & 2 deletions projects/rccl/src/algorithms/gin/sdma/gin_alltoall_sdma.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand All @@ -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;
Comment thread
Kapil-Shyam-Pawar marked this conversation as resolved.
Resolved

// Every rank must be reachable over LSA.
if (ncclTeamLsa(comm).nRanks != comm->nRanks) return false;
Expand Down
9 changes: 9 additions & 0 deletions projects/rccl/src/include/nccl_device/gin/anvil_sdma/gin_anvil_sdma.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -469,4 +469,13 @@ struct ncclGinApi_Wait<NCCL_NET_DEVICE_GIN_ANVIL_SDMA> {
}
};

// Only Flush drains the SDMA engine, so the barrier has to keep its pre-signal
// fenceFlush.
template <>
struct ncclGinApi_SupportsStrongSignal<NCCL_NET_DEVICE_GIN_ANVIL_SDMA> {
NCCL_DEVICE_INLINE static bool call(ncclGinCtx) {
return false;
}
};

#endif
8 changes: 8 additions & 0 deletions projects/rccl/src/include/nccl_device/gin/rocshmem_gda/gin_rocshmem_gda.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -186,4 +186,12 @@ struct ncclGinApi_Wait<NCCL_NET_DEVICE_GIN_ROCSHMEM_GDA> {
}
};

// Mirrors ginRocshmemGdaGetGinProperties()'s supportsStrongSignals on the host side.
template <>
struct ncclGinApi_SupportsStrongSignal<NCCL_NET_DEVICE_GIN_ROCSHMEM_GDA> {
NCCL_DEVICE_INLINE static bool call(ncclGinCtx) {
return true;
}
};

#endif /* _NCCL_DEVICE_GIN_ROCSHMEM_GDA_H_ */
19 changes: 15 additions & 4 deletions projects/rccl/test/GinAlltoAllEligibilityTests.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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; }
Expand Down Expand Up @@ -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;
Comment thread
Kapil-Shyam-Pawar marked this conversation as resolved.
Resolved
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.
Expand Down
4 changes: 2 additions & 2 deletions projects/rccl/tools/ci/lib/gin-tests.json
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"],
Comment thread
Kapil-Shyam-Pawar marked this conversation as resolved.
Resolved
"args": "-b 128 -e {E} -f 2 -g 1 -R 2 -D 3 -A 1"
},
{
Expand Down Expand Up @@ -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"
}
]
Expand Down
4 changes: 2 additions & 2 deletions projects/rccl/tools/ci/run-gin-ci.sh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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}" \
Expand Down
12 changes: 6 additions & 6 deletions projects/rccl/tools/scripts/test_runner/configs/gin_rocshmem_backends_mellanox.json
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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" },
Expand All @@ -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" },
Expand All @@ -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" },
Expand All @@ -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" }
Expand Down Expand Up @@ -191,7 +191,7 @@
"gin_devmpi_gda": {
"extends": "gin_devmpi_common",
"env_variables": {
"NCCL_GIN_TYPE": "5"
"NCCL_GIN_TYPE": "6"
}
},
"gin_devmpi_sdma": {
Expand All @@ -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 },
Expand Down
Loading