Skip to content

fix(gin): de-alias Anvil signals across logical contexts #11569

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

Closed
dlamd1dai wants to merge 8 commits into from
Closed

fix(gin): de-alias Anvil signals across logical contexts #11569

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
8 commits
Select commit Hold shift + click to select a range
42d25db
fix(gin): de-alias Anvil signals across contexts
dlamd1dai Sep 14, 2026
d3bbb0c
fix(gin): address ARGUS review and GIN-CI poison failures for Anvil c…
dlamd1dai Sep 14, 2026
041d7b0
test(gin): cover PutValue logical contextId; split rocSHMEM atomics
dlamd1dai Sep 14, 2026
34deac8
fix(gin): compile Anvil bind-signals and MPI stripe assertions
dlamd1dai Sep 14, 2026
7325296
fix(gin): address post-rebase ARGUS findings for Anvil contexts
dlamd1dai Sep 29, 2026
802b251
fix(gin): index Anvil signal peers in GIN-team space
dlamd1dai Sep 29, 2026
59994f4
fix(gin): run Anvil bind collectives once per signal span
dlamd1dai Sep 29, 2026
cf8cf97
fix(gin): size Anvil signal allgather to world; lock the pending walk
dlamd1dai Oct 1, 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
397 changes: 263 additions & 134 deletions projects/rccl/src/gin/gin_plugin_anvil_sdma.cc
View file
Open in desktop

Large diffs are not rendered by default.

2 changes: 1 addition & 1 deletion projects/rccl/src/include/gin/gin_host_anvil_sdma.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -15,7 +15,7 @@ extern ncclGin_t ncclGinAnvilSdmaPlugin;

struct ncclComm;
ncclResult_t ncclGinAnvilBindResourceWindowSignals(struct ncclComm* comm, void* resourceUserPtr, size_t arenaByteOffset,
int nContexts, int nSignalsPerContext);
int nSignalSlots, int nSignalsPerContext);

void ncclGinAnvilSetInitContext(void* initCtx, struct ncclComm* comm);

Expand Down
24 changes: 15 additions & 9 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 @@ -35,6 +35,12 @@ NCCL_DEVICE_INLINE bool anvilCtxValid(ncclGinAnvilSdmaGPUContext* rsCtx) {
return rsCtx != nullptr && loadConst(&rsCtx->layoutMagic) == NCCL_GIN_ANVIL_SDMA_LAYOUT_MAGIC;
}

// Logical contexts index distinct GPU-context slots. Counters are sliced per
// context (countersDev + contextId * nCounters), matching proxy/gdaki.
NCCL_DEVICE_INLINE ncclGinAnvilSdmaGPUContext* anvilGpuCtx(ncclGinCtx ctx) {
Comment thread
dlamd1dai marked this conversation as resolved.
Resolved
return &reinterpret_cast<ncclGinAnvilSdmaGPUContext*>(ctx.handle)[ctx.contextId];
}

NCCL_DEVICE_INLINE void* resolveRemotePeerVa(ncclGinAnvilSdmaGPUContext* rsCtx, ncclGinAnvilSdmaMemHandle* mh, int peer,
size_t off) {
ptrdiff_t stride = loadConst(&mh->vmmStride);
Expand Down Expand Up @@ -183,7 +189,7 @@ struct ncclGinApi_Put<NCCL_NET_DEVICE_GIN_ANVIL_SDMA> {

if (coop.thread_rank() != 0) return;

ncclGinAnvilSdmaGPUContext* rsCtx = (ncclGinAnvilSdmaGPUContext*)ctx.handle;
ncclGinAnvilSdmaGPUContext* rsCtx = nccl::gin::anvil::detail::anvilGpuCtx(ctx);
if (!anvilCtxValid(rsCtx)) return;
const int blockId = blockIdx.x + blockIdx.y * gridDim.x;

Expand Down Expand Up @@ -298,7 +304,7 @@ struct ncclGinApi_PutValue<NCCL_NET_DEVICE_GIN_ANVIL_SDMA> {

if (coop.thread_rank() != 0) return;

ncclGinAnvilSdmaGPUContext* rsCtx = (ncclGinAnvilSdmaGPUContext*)ctx.handle;
ncclGinAnvilSdmaGPUContext* rsCtx = nccl::gin::anvil::detail::anvilGpuCtx(ctx);
if (!anvilCtxValid(rsCtx)) return;
const int blockId = blockIdx.x + blockIdx.y * gridDim.x;
ncclGinAnvilSdmaMemHandle* dstMh = (ncclGinAnvilSdmaMemHandle*)dstWin;
Expand Down Expand Up @@ -367,7 +373,7 @@ struct ncclGinApi_PutValue<NCCL_NET_DEVICE_GIN_ANVIL_SDMA> {
template <>
struct ncclGinApi_GetCounterPtr<NCCL_NET_DEVICE_GIN_ANVIL_SDMA> {
NCCL_DEVICE_INLINE static ncclGinOffsetPtr call(ncclGinCtx ctx, ncclGinCounter_t counterId) {
ncclGinAnvilSdmaGPUContext* rsCtx = (ncclGinAnvilSdmaGPUContext*)ctx.handle;
ncclGinAnvilSdmaGPUContext* rsCtx = nccl::gin::anvil::detail::anvilGpuCtx(ctx);
if (!nccl::gin::anvil::detail::anvilCtxValid(rsCtx)) return {nullptr, 0};
return {nccl::utility::loadConst(&rsCtx->counters) + counterId, 0};
}
Expand All @@ -376,7 +382,7 @@ struct ncclGinApi_GetCounterPtr<NCCL_NET_DEVICE_GIN_ANVIL_SDMA> {
template <>
struct ncclGinApi_ResetCounter<NCCL_NET_DEVICE_GIN_ANVIL_SDMA> {
NCCL_DEVICE_INLINE static void call(ncclGinCtx ctx, ncclGinCounter_t counterId) {
ncclGinAnvilSdmaGPUContext* rsCtx = (ncclGinAnvilSdmaGPUContext*)ctx.handle;
ncclGinAnvilSdmaGPUContext* rsCtx = nccl::gin::anvil::detail::anvilGpuCtx(ctx);
if (!nccl::gin::anvil::detail::anvilCtxValid(rsCtx)) return;
uint64_t* counters = nccl::utility::loadConst(&rsCtx->counters);
if (counters != nullptr) counters[counterId] = 0;
Expand All @@ -388,7 +394,7 @@ struct ncclGinApi_GetSignalPtr<NCCL_NET_DEVICE_GIN_ANVIL_SDMA> {
NCCL_DEVICE_INLINE static ncclGinOffsetPtr call(ncclGinCtx ctx, ncclGinSignal_t signalId) {
using nccl::gin::anvil::detail::anvilCtxValid;
using nccl::utility::loadConst;
ncclGinAnvilSdmaGPUContext* rsCtx = (ncclGinAnvilSdmaGPUContext*)ctx.handle;
ncclGinAnvilSdmaGPUContext* rsCtx = nccl::gin::anvil::detail::anvilGpuCtx(ctx);
assert(anvilCtxValid(rsCtx));
uint64_t* signals = loadConst(&rsCtx->signals);
assert(signals != nullptr);
Expand All @@ -399,7 +405,7 @@ struct ncclGinApi_GetSignalPtr<NCCL_NET_DEVICE_GIN_ANVIL_SDMA> {
template <>
struct ncclGinApi_ResetSignal<NCCL_NET_DEVICE_GIN_ANVIL_SDMA> {
NCCL_DEVICE_INLINE static void call(ncclGinCtx ctx, ncclGinSignalDescriptor signal) {
ncclGinAnvilSdmaGPUContext* rsCtx = (ncclGinAnvilSdmaGPUContext*)ctx.handle;
ncclGinAnvilSdmaGPUContext* rsCtx = nccl::gin::anvil::detail::anvilGpuCtx(ctx);
if (!nccl::gin::anvil::detail::anvilCtxValid(rsCtx)) return;
if (signal.type == NCCL_GIN_SIGNAL_TYPE_INDEXED) {
uint64_t* signals = nccl::utility::loadConst(&rsCtx->signals);
Expand All @@ -418,7 +424,7 @@ struct ncclGinApi_Flush<NCCL_NET_DEVICE_GIN_ANVIL_SDMA> {
(void)ord;
(void)abortFlag;
using nccl::utility::loadConst;
ncclGinAnvilSdmaGPUContext* rsCtx = (ncclGinAnvilSdmaGPUContext*)ctx.handle;
ncclGinAnvilSdmaGPUContext* rsCtx = nccl::gin::anvil::detail::anvilGpuCtx(ctx);
if (!nccl::gin::anvil::detail::anvilCtxValid(rsCtx)) {
__threadfence_system();
return;
Expand Down Expand Up @@ -475,7 +481,7 @@ struct ncclGinApi_Get<NCCL_NET_DEVICE_GIN_ANVIL_SDMA> {

if (coop.thread_rank() != 0 || bytes == 0) return;

ncclGinAnvilSdmaGPUContext* rsCtx = (ncclGinAnvilSdmaGPUContext*)ctx.handle;
ncclGinAnvilSdmaGPUContext* rsCtx = nccl::gin::anvil::detail::anvilGpuCtx(ctx);
if (!anvilCtxValid(rsCtx)) return;
const int blockId = blockIdx.x + blockIdx.y * gridDim.x;
ncclGinAnvilSdmaMemHandle* remoteMh = (ncclGinAnvilSdmaMemHandle*)remoteWin;
Expand Down Expand Up @@ -519,7 +525,7 @@ struct ncclGinApi_FlushAsync<NCCL_NET_DEVICE_GIN_ANVIL_SDMA> {
auto* request = reinterpret_cast<nccl::gin::anvil::detail::ncclGinAnvilSdmaRequest*>(outRequest);
request->complete = 0;

ncclGinAnvilSdmaGPUContext* rsCtx = (ncclGinAnvilSdmaGPUContext*)ctx.handle;
ncclGinAnvilSdmaGPUContext* rsCtx = nccl::gin::anvil::detail::anvilGpuCtx(ctx);
if (nccl::gin::anvil::detail::anvilCtxValid(rsCtx)) {
uint64_t* dirtyPtr = loadConst(&rsCtx->sdmaDirty);
uint64_t dirty =
Expand Down
112 changes: 112 additions & 0 deletions projects/rccl/test/device/GinAnvilSdmaTemplate_test.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -396,6 +396,118 @@ TEST_F(GinAnvilSdmaTemplateTest, CounterSignal_GetReset) {
EXPECT_EQ(d_outSig.download(), 0ULL);
}

// AICOMRCCL-2339: the Anvil device dispatch must select the GPU context
// indexed by ncclGinCtx::contextId instead of always using array element zero.
__global__ void kernelSignalContextSelection(ncclGinAnvilSdmaGPUContext* contexts) {
if (threadIdx.x != 0) return;
ncclGinCtx ginCtx{};
ginCtx.handle = contexts;

ginCtx.contextId = 0;
ncclGinApi_GetSignalPtr<NCCL_NET_DEVICE_GIN_ANVIL_SDMA>::call(ginCtx, 0).ptr[0] = 11;
ncclGinApi_GetCounterPtr<NCCL_NET_DEVICE_GIN_ANVIL_SDMA>::call(ginCtx, 0).ptr[0] = 101;

ginCtx.contextId = 1;
ncclGinApi_GetSignalPtr<NCCL_NET_DEVICE_GIN_ANVIL_SDMA>::call(ginCtx, 0).ptr[0] = 22;
ncclGinApi_ResetSignal<NCCL_NET_DEVICE_GIN_ANVIL_SDMA>::call(
ginCtx, ncclGinSignalDescriptor{NCCL_GIN_SIGNAL_TYPE_INDEXED, {.indexedSignal = {.signalId = 0}}});
ncclGinApi_ResetCounter<NCCL_NET_DEVICE_GIN_ANVIL_SDMA>::call(ginCtx, 0);
ncclGinApi_GetCounterPtr<NCCL_NET_DEVICE_GIN_ANVIL_SDMA>::call(ginCtx, 0).ptr[0] = 202;
}

TEST_F(GinAnvilSdmaTemplateTest, SignalApis_SelectLogicalContext) {
Comment thread
dlamd1dai marked this conversation as resolved.
Resolved
DeviceBuffer<uint64_t> d_signals(2);
DeviceBuffer<uint64_t> d_counters(2);
d_signals.zero();
d_counters.zero();
ncclGinAnvilSdmaGPUContext hostCtx[2]{};
for (int i = 0; i < 2; i++) {
hostCtx[i].layoutMagic = NCCL_GIN_ANVIL_SDMA_LAYOUT_MAGIC;
hostCtx[i].signals = d_signals.ptr + i;
hostCtx[i].counters = d_counters.ptr + i;
hostCtx[i].nSignals = 1;
hostCtx[i].nCounters = 1;
}
DeviceBuffer<ncclGinAnvilSdmaGPUContext> d_contexts(2);
d_contexts.copyFrom(hostCtx, 2);

kernelSignalContextSelection<<<1, 1>>>(d_contexts.ptr);
syncAndCheck();
auto signals = d_signals.copyTo();
auto counters = d_counters.copyTo();
EXPECT_EQ(signals[0], 11ULL) << "ResetSignal must not clear context 0 when invoked on contextId=1";
EXPECT_EQ(signals[1], 0ULL) << "ResetSignal on contextId=1 must clear context 1 signal cell";
EXPECT_EQ(counters[0], 101ULL);
EXPECT_EQ(counters[1], 202ULL);
}

// AICOMRCCL-2339: PutValue must honor ncclGinCtx::contextId when resolving signals.
__global__ void kernelPutValueContextSelection(ncclGinAnvilSdmaGPUContext* contexts, ncclGinAnvilSdmaMemHandle* dstMh) {
if (threadIdx.x != 0) return;
ncclGinCtx ginCtx{};
ginCtx.handle = contexts;
ginCtx.nRanks = 2;
ginCtx.rank = 0;
ginCtx.contextId = 1;

ncclGinSignalDescriptor sig{};
sig.type = NCCL_GIN_SIGNAL_TYPE_INDEXED;
sig.indexedSignal.signalId = 0;
ncclGinApi_PutValue<NCCL_NET_DEVICE_GIN_ANVIL_SDMA>::call(
ginCtx, ncclCoopThread{}, 1, reinterpret_cast<ncclGinWindow_t>(dstMh), 0, static_cast<uint8_t>(0x11), sig,
ncclGinSignalInc, 1, false, nullptr, cuda::thread_scope_system, cuda::thread_scope_system);
}

TEST_F(GinAnvilSdmaTemplateTest, PutValue_SelectLogicalContext) {
DeviceBuffer<uint64_t> d_signals(2);
DeviceBuffer<uint8_t> d_dst(8);
DeviceBuffer<uintptr_t> d_remoteAddrs0(2);
DeviceBuffer<uintptr_t> d_remoteAddrs1(2);
d_signals.zero();
d_dst.zero();

// Distinct peer-1 fallback entries so reverting anvilGpuCtx to handle[0]
// would write context 0's cell and fail the assertions below.
uintptr_t remoteAddrs0[2] = {0, reinterpret_cast<uintptr_t>(d_signals.ptr + 0)};
uintptr_t remoteAddrs1[2] = {0, reinterpret_cast<uintptr_t>(d_signals.ptr + 1)};
d_remoteAddrs0.copyFrom(remoteAddrs0, 2);
d_remoteAddrs1.copyFrom(remoteAddrs1, 2);

DeviceBuffer<ncclGinAnvilIpcBufEntry> d_entry(1);
ncclGinAnvilIpcBufEntry ipcEntry{};
ipcEntry.local_base = reinterpret_cast<uintptr_t>(d_dst.ptr);
ipcEntry.length = 4096;
ipcEntry.remote_bases[1] = reinterpret_cast<uintptr_t>(d_dst.ptr);
d_entry.upload(ipcEntry);

ncclGinAnvilSdmaMemHandle dstMh{};
dstMh.baseAddr = reinterpret_cast<uintptr_t>(d_dst.ptr);
DeviceBuffer<ncclGinAnvilSdmaMemHandle> d_dstMh(1);
d_dstMh.upload(dstMh);

ncclGinAnvilSdmaGPUContext hostCtx[2]{};
for (int i = 0; i < 2; i++) {
hostCtx[i].layoutMagic = NCCL_GIN_ANVIL_SDMA_LAYOUT_MAGIC;
hostCtx[i].signals = d_signals.ptr + i;
hostCtx[i].signal_remote_addrs = (i == 0) ? d_remoteAddrs0.ptr : d_remoteAddrs1.ptr;
hostCtx[i].ipcTable = d_entry.ptr;
hostCtx[i].ipcTableCount = 1;
hostCtx[i].nRanks = 2;
hostCtx[i].rank = 0;
hostCtx[i].sdmaThreshold = 4096;
hostCtx[i].fusedSdmaSignal = 0;
}

DeviceBuffer<ncclGinAnvilSdmaGPUContext> d_contexts(2);
d_contexts.copyFrom(hostCtx, 2);

kernelPutValueContextSelection<<<1, 1>>>(d_contexts.ptr, d_dstMh.ptr);
syncAndCheck();
auto signals = d_signals.copyTo();
EXPECT_EQ(signals[0], 0ULL) << "context 0 signal stripe must stay untouched";
EXPECT_EQ(signals[1], 1ULL) << "PutValue on contextId=1 must increment context 1 signal cell";
}

// H10: invalid ctx on getters returns nullptr / no-op.
__global__ void kernelInvalidCtxApis(bool* ok) {
ncclGinCtx ginCtx{};
Expand Down
138 changes: 138 additions & 0 deletions projects/rccl/test/gin/GinAnvilPlugin_test.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -364,6 +364,65 @@ TEST_F(GinAnvilPluginTest, BindSignals_Success) {
plugin_.finalize(ictx);
}

// AICOMRCCL-2339: logical contexts on one Anvil connection share SDMA
// queues, but each context must bind a distinct signal stripe. A second
// pending context on the same comm makes slot and contextId diverge: the
// first context's contextId 0 is not slot 0.
TEST_F(GinAnvilPluginTest, BindSignals_LogicalContextsUseDistinctSignalStripes) {
Comment thread
dlamd1dai marked this conversation as resolved.
Resolved
void* ictx = nullptr;
initCtx(&ictx);
void* coll = nullptr;
connectColl(ictx, &coll);
ncclGinConfig_t cfg{};
cfg.nContexts = 3;
cfg.nSignals = 2;
cfg.nCounters = 5;
void* ginCtx = nullptr;
ncclNetDeviceHandle_v11_t* devHandle = nullptr;
ASSERT_EQ(plugin_.createContext(coll, &cfg, &ginCtx, &devHandle), ncclSuccess);
Comment thread
dlamd1dai marked this conversation as resolved.
Resolved

ncclGinConfig_t cfgB{};
cfgB.nContexts = 1;
cfgB.nSignals = 2;
cfgB.nCounters = 5;
void* ginCtxB = nullptr;
ncclNetDeviceHandle_v11_t* devHandleB = nullptr;
ASSERT_EQ(plugin_.createContext(coll, &cfgB, &ginCtxB, &devHandleB), ncclSuccess);

char arena[8192] = {};
// Pending list is newest-first: ginCtxB takes slot 0, ginCtx takes slots 1..3.
ASSERT_EQ(ncclGinAnvilBindResourceWindowSignals(mockComm_.get(), arena, 0, 4, 4), ncclSuccess);
ASSERT_EQ(devHandle->size, 3 * sizeof(ncclGinAnvilSdmaGPUContext));
ASSERT_EQ(devHandleB->size, sizeof(ncclGinAnvilSdmaGPUContext));

ncclGinAnvilSdmaGPUContext hostCtx[3]{};
ncclGinAnvilSdmaGPUContext hostCtxB{};
ASSERT_EQ(hipMemcpy(hostCtx, devHandle->handle, sizeof(hostCtx), hipMemcpyDeviceToHost), hipSuccess);
ASSERT_EQ(hipMemcpy(&hostCtxB, devHandleB->handle, sizeof(hostCtxB), hipMemcpyDeviceToHost), hipSuccess);
ASSERT_NE(hostCtx[0].signals, nullptr);
ASSERT_NE(hostCtx[1].signals, nullptr);
ASSERT_NE(hostCtx[2].signals, nullptr);
ASSERT_NE(hostCtxB.signals, nullptr);
EXPECT_EQ(hostCtx[1].signals - hostCtx[0].signals, 4);
EXPECT_EQ(hostCtx[2].signals - hostCtx[1].signals, 4);
EXPECT_EQ(hostCtx[0].signals - hostCtxB.signals, 4);
EXPECT_NE(hostCtx[0].signal_remote_addrs, hostCtx[1].signal_remote_addrs);

// Queue ownership remains at the connection level by design.
EXPECT_EQ(hostCtx[0].queueHandles, hostCtx[1].queueHandles);
EXPECT_EQ(hostCtx[0].sdmaDirty, hostCtx[1].sdmaDirty);
EXPECT_EQ(hostCtx[0].queueHandles, hostCtxB.queueHandles);
EXPECT_NE(hostCtx[0].counters, nullptr);
EXPECT_NE(hostCtx[1].counters, nullptr);
EXPECT_EQ(hostCtx[1].counters - hostCtx[0].counters, 5);
EXPECT_EQ(hostCtx[2].counters - hostCtx[1].counters, 5);

plugin_.destroyContext(ginCtx);
plugin_.destroyContext(ginCtxB);
plugin_.closeColl(coll);
plugin_.finalize(ictx);
}

// G15: dereg null, progress, queryLastError.
TEST_F(GinAnvilPluginTest, Misc_NoOpPaths) {
EXPECT_EQ(plugin_.deregMrSym(nullptr, nullptr), ncclSuccess);
Expand Down Expand Up @@ -661,6 +720,85 @@ TEST_F(GinAnvilPluginTest, ConnCheck_HealthyConnectivitySucceeds) {
plugin_.finalize(ictx);
}

// AICOMRCCL-2339: the comm-wide collectives in the bind path are per span, not
// per logical context. Binding a context with nContexts=3 must still cost one
// local-base allgather and one conn-check, or ncclDevCommCreate pays a bootstrap
// round trip for every logical context.
TEST_F(GinAnvilPluginTest, BindSignals_CommWideCollectivesRunOncePerSpan) {
void* rawDevLsa = nullptr;
ASSERT_EQ(hipMalloc(&rawDevLsa, sizeof(uint64_t) * 8), hipSuccess);
HipAllocation devLsa(rawDevLsa);
GinAnvilPluginStubs::SetLsaSelfAddr(devLsa.get());
GinAnvilPluginStubs::SetBootstrapNranks(2);
mockComm_.get()->devrState.lsaSize = 2;

void* ictx = nullptr;
initCtx(&ictx);
void* coll = nullptr;
connectColl(ictx, &coll, 2);
ncclGinConfig_t cfg{};
cfg.nContexts = 3;
cfg.nSignals = 2;
void* ginCtx = nullptr;
ncclNetDeviceHandle_v11_t* devHandle = nullptr;
ASSERT_EQ(plugin_.createContext(coll, &cfg, &ginCtx, &devHandle), ncclSuccess);

// connectColl runs the factory allgather; only the bind path is under test.
const int allGathersBeforeBind = GinAnvilPluginStubs::GetBootstrapAllGatherCalls();
const int intraNodeBeforeBind = GinAnvilPluginStubs::GetIntraNodeAllGatherCalls();

char arena[8192] = {};
ASSERT_EQ(ncclGinAnvilBindResourceWindowSignals(mockComm_.get(), arena, 0, 3, 2), ncclSuccess);

EXPECT_EQ(GinAnvilPluginStubs::GetBootstrapAllGatherCalls() - allGathersBeforeBind, 1);
// One conn-check, which costs two LSA allgathers: the setup state exchange and
// the missing-count exchange of its single attempt.
EXPECT_EQ(GinAnvilPluginStubs::GetIntraNodeAllGatherCalls() - intraNodeBeforeBind, 2);
EXPECT_EQ(GinAnvilPluginStubs::GetConnCheckWriteCalls(), 1);
EXPECT_EQ(GinAnvilPluginStubs::GetConnCheckVerifyCalls(), 1);

// All three contexts are still bound to their own stripe.
ncclGinAnvilSdmaGPUContext hostCtx[3]{};
ASSERT_EQ(hipMemcpy(hostCtx, devHandle->handle, sizeof(hostCtx), hipMemcpyDeviceToHost), hipSuccess);
EXPECT_EQ(hostCtx[1].signals - hostCtx[0].signals, 2);
EXPECT_EQ(hostCtx[2].signals - hostCtx[1].signals, 2);

plugin_.destroyContext(ginCtx);
plugin_.closeColl(coll);
plugin_.finalize(ictx);
}

// The skip path returns before the comm is marked as checked, so a per-context
// conn-check call would repeat its setup allgather nContexts times.
TEST_F(GinAnvilPluginTest, BindSignals_SkippedConnCheckStillRunsOncePerSpan) {
GinAnvilPluginStubs::SetBootstrapNranks(2);
mockComm_.get()->devrState.lsaSize = 4; // != nRanks, so the gate skips

void* ictx = nullptr;
initCtx(&ictx);
void* coll = nullptr;
connectColl(ictx, &coll, 2);
ncclGinConfig_t cfg{};
cfg.nContexts = 3;
cfg.nSignals = 2;
void* ginCtx = nullptr;
ncclNetDeviceHandle_v11_t* devHandle = nullptr;
ASSERT_EQ(plugin_.createContext(coll, &cfg, &ginCtx, &devHandle), ncclSuccess);

const int intraNodeBeforeBind = GinAnvilPluginStubs::GetIntraNodeAllGatherCalls();

char arena[8192] = {};
ASSERT_EQ(ncclGinAnvilBindResourceWindowSignals(mockComm_.get(), arena, 0, 3, 2), ncclSuccess);

EXPECT_EQ(GinAnvilPluginStubs::GetIntraNodeAllGatherCalls() - intraNodeBeforeBind, 1);
EXPECT_EQ(GinAnvilPluginStubs::GetConnCheckWriteCalls(), 0);
EXPECT_EQ(GinAnvilPluginStubs::GetConnCheckVerifyCalls(), 0);

plugin_.destroyContext(ginCtx);
plugin_.closeColl(coll);
plugin_.finalize(ictx);
}

TEST_F(GinAnvilPluginTest, ConnCheck_NonNumericEnvValueKeepsGateEnabled) {
ScopedEnv enabled("NCCL_GIN_ANVIL_SDMA_CONN_CHECK", "true");
void* rawDevLsa = nullptr;
Expand Down
Loading
Loading