diff --git a/projects/rccl/CHANGELOG.md b/projects/rccl/CHANGELOG.md index e80b3c32ae78..568ffe7cbc63 100644 --- a/projects/rccl/CHANGELOG.md +++ b/projects/rccl/CHANGELOG.md @@ -56,7 +56,7 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https: * Fixed virtual address space exhaustion when symmetric windows backed by the same physical allocation are registered repeatedly. * Fixed Copy Engine `ncclAlltoAllv` hanging on intra-node communicators of more than 8 ranks with uneven send sizes. Each rank decided from its own send sizes whether to add intra-batch barriers, so ranks could issue different barrier counts. The decision now uses the full send-size matrix, which every rank shares. * Fixed `NCCL_P2P_DISABLE` and `NCCL_P2P_LEVEL` being ignored on Intel, ARM and Zhaoxin hosts, where the CPU-based default P2P level overwrote the user setting. On such hosts `NCCL_P2P_DISABLE=1` now falls back to SHM as documented, and `NCCL_P2P_LEVEL` is honored. -* Fixed a 15 to 30% `ncclScatter` bandwidth loss from 8 MB upward on multi-node AINIC jobs of up to 8 nodes, caused by the move to two network channels per peer. Scatter now uses and connects one channel per peer when the peers cannot each get their own channels. Point-to-point operations and the other collectives keep two channels. +* Fixed a 15 to 30% `ncclScatter` bandwidth loss from 16 MB upward on multi-node AINIC jobs, caused by the move to two network channels per peer (for example on 4 nodes of 8 GPUs). Scatter now uses and connects the most channels per peer whose flows still fit the P2P channel pool. Point-to-point operations and the other collectives keep their channel count. ### Known issues * The FP8 ReduceSum and ReduceCopy device APIs are not exported in the LLVM bitcode library. diff --git a/projects/rccl/src/enqueue/enqueue.cc b/projects/rccl/src/enqueue/enqueue.cc index cfb2b22d8477..aa3e388a22f3 100644 --- a/projects/rccl/src/enqueue/enqueue.cc +++ b/projects/rccl/src/enqueue/enqueue.cc @@ -1326,19 +1326,16 @@ static bool rcclP2pBatchEligible(struct ncclComm* comm, ssize_t sendBytes, ssize rcclParamP2pBatchThreshold()); } -// Number of the p2pnChannelsPerPeer parts that a task of collAPI uses and connects. -// On AINIC ncclTopoComputeP2pChannels raises p2pnChannelsPerPeer to 2 so one peer can saturate a NIC. -// Scatter already drives all nRanks-1 peers at once from one root and loses 15-30% from 8 MB up with -// two parts per peer, so it uses only part 0 once the peers cannot each own p2pnChannelsPerPeer -// channels. It must not connect part 1 either: when scatter connects both parts itself, it stays at -// or below the two-part speed even though it sends on part 0 only. Part 1 connected later by another -// collective does not slow it down. The answer may depend only on values both ends of a pair share -// (collAPI, comm-wide sizes), never on per-rank task counts, or sender and receiver would disagree on -// the parts. +// Number of the p2pnChannelsPerPeer parts a task of collAPI uses and connects. Multi-node AINIC scatter +// uses and connects only the most parts whose flows from the root still fit the pool; connecting more +// slows it down. A user-set nChannelsPerNetPeer is kept, as in ncclTopoComputeP2pChannels. Depend only +// on values both ends of a pair share, or they disagree on the parts. int rcclP2pTaskParts(struct ncclComm* comm, ncclFunc_t collAPI) { - if (collAPI == ncclFuncScatter && comm->nNodes > 1 && comm->p2pnChannelsPerPeer > 1 && rcclUseAinic() && - (comm->nRanks - 1) * comm->p2pnChannelsPerPeer > comm->p2pnChannels) - return 1; + if (collAPI == ncclFuncScatter && comm->nNodes > 1 && comm->config.nChannelsPerNetPeer == NCCL_CONFIG_UNDEF_INT && + rcclUseAinic()) { + int partsThatFit = comm->p2pnChannels / (comm->nRanks - 1); + return std::min(comm->p2pnChannelsPerPeer, partsThatFit >= 1 ? pow2Down(partsThatFit) : 1); + } return comm->p2pnChannelsPerPeer; } @@ -1348,16 +1345,17 @@ static int rcclP2pPolicyChannels(struct ncclComm* comm, struct ncclTaskP2p* task ncclFunc_t collAPI; int nChannels; }; - // This workaround targets direct P2P/IPC traffic only. When P2P is disabled, - // preserve the normal channel selection for the fallback transport. if (task == nullptr) return -1; + // The Scatter part limit is about network flows, so it applies with P2P disabled too. int taskParts = rcclP2pTaskParts(comm, task->collAPI); if (taskParts < comm->p2pnChannelsPerPeer) return taskParts; + // The policies below target direct P2P/IPC traffic only. When P2P is disabled, + // preserve the normal channel selection for the fallback transport. if (ncclParamP2pDisable()) return -1; constexpr Policy policies[] = { // Two-channel full-mesh traffic has high variance on gfx110x. Keep the - // workaround specific to AllToAll so Gather, Scatter, and SendRecv retain - // their higher-throughput multi-channel paths. + // workaround specific to AllToAll so Gather, SendRecv and Scatter (outside the + // AINIC part limit above) retain their higher-throughput multi-channel paths. {"gfx110", ncclFuncAlltoAll, 1}, }; for (const auto& policy : policies) { diff --git a/projects/rccl/src/enqueue/task_prep/task_posttuning.cc b/projects/rccl/src/enqueue/task_prep/task_posttuning.cc index 216aab7e8696..527840d47ff6 100644 --- a/projects/rccl/src/enqueue/task_prep/task_posttuning.cc +++ b/projects/rccl/src/enqueue/task_prep/task_posttuning.cc @@ -245,8 +245,7 @@ static ncclResult_t postTuneP2pRegisterBuffer(struct ncclComm* comm, struct nccl if (network) { bool pxnUsed = !ncclPxnDisable(comm) && comm->isAllNvlink && comm->maxLocalRanks > 1; if (proxySameProcess && !pxnUsed && (conn->conn.flags & NCCL_DIRECT_NIC)) { - int nParts = rcclP2pTaskParts(comm, task->collAPI); - for (int part = 0; part < nParts; part++) { + for (int part = 0; part < comm->p2pnChannelsPerPeer; part++) { int partChannelId = ncclP2pChannelForPart(comm->p2pnChannels, base, part, comm->p2pnChannelsPerPeer, comm->nNodes, comm->p2pChannelShiftSize); struct ncclConnector* partConn = isSendNotRecv ? &comm->channels[partChannelId].peers[peer]->send[connIndex] : diff --git a/projects/rccl/test/host/enqueue-test.cc b/projects/rccl/test/host/enqueue-test.cc index 6830f01a4afe..6374f2dcb457 100644 --- a/projects/rccl/test/host/enqueue-test.cc +++ b/projects/rccl/test/host/enqueue-test.cc @@ -339,6 +339,68 @@ TEST_F(EnqueueMicrotest, ShmemDynamicSize_DeviceLinker_IsAlwaysZero) { } #endif +// =========================================================================== +// rcclP2pTaskParts (enqueue.cc:1333) and rcclP2pPolicyChannels (enqueue.cc:1342) +// +// Multi-node AINIC Scatter takes the most power-of-two parts whose nRanks - 1 +// flows fit p2pnChannels, at most p2pnChannelsPerPeer; everything else takes +// p2pnChannelsPerPeer. Both ends of a pair must get the same answer, so it may +// depend only on comm-wide fields. +// =========================================================================== + +namespace { +std::unique_ptr P2pPartsComm(int nNodes, int nRanks, int p2pnChannels, int perPeer) { + std::unique_ptr comm{new ncclComm{}}; + comm->nNodes = nNodes; + comm->nRanks = nRanks; + comm->p2pnChannels = p2pnChannels; + comm->p2pnChannelsPerPeer = perPeer; + comm->config.nChannelsPerNetPeer = NCCL_CONFIG_UNDEF_INT; + return comm; +} +} // namespace + +TEST_F(EnqueueMicrotest, P2pTaskParts_AinicScatterTakesThePartsWhoseFlowsFitThePool) { + g_rcclUseAinicValue = true; + struct Case { + int nNodes, nRanks, p2pnChannels, perPeer, scatterParts; + }; + // 8 GPUs per node with the MI355X pools, the pool boundary (16 peers x 2 parts fill 32 channels, + // 17 peers do not), power-of-two rounding (64 / 19 = 3) and the one-part floor (64 / 127 = 0). + for (Case k : {Case{2, 16, 32, 2, 2}, Case{3, 24, 64, 4, 2}, Case{4, 32, 32, 2, 1}, Case{8, 64, 64, 2, 1}, + Case{2, 17, 32, 2, 2}, Case{2, 18, 32, 2, 1}, Case{3, 20, 64, 8, 2}, Case{16, 128, 64, 2, 1}}) { + auto comm = P2pPartsComm(k.nNodes, k.nRanks, k.p2pnChannels, k.perPeer); + EXPECT_EQ(k.scatterParts, rcclP2pTaskParts(comm.get(), ncclFuncScatter)) + << k.nRanks << " ranks, " << k.p2pnChannels << " channels"; + for (ncclFunc_t f : {ncclFuncSendRecv, ncclFuncSend, ncclFuncRecv, ncclFuncAlltoAll, ncclFuncGather, + ncclFuncAlltoAllv}) { + EXPECT_EQ(k.perPeer, rcclP2pTaskParts(comm.get(), f)) << "collAPI=" << int(f); + } + } +} + +TEST_F(EnqueueMicrotest, P2pTaskParts_ScatterKeepsPerPeerOutsideMultiNodeAinicOrWithAUserCount) { + auto comm = P2pPartsComm(/*nNodes=*/4, /*nRanks=*/32, /*p2pnChannels=*/32, /*perPeer=*/2); + g_rcclUseAinicValue = false; + EXPECT_EQ(2, rcclP2pTaskParts(comm.get(), ncclFuncScatter)) << "not AINIC"; + g_rcclUseAinicValue = true; + comm->config.nChannelsPerNetPeer = 2; + EXPECT_EQ(2, rcclP2pTaskParts(comm.get(), ncclFuncScatter)) << "NCCL_NCHANNELS_PER_NET_PEER set"; + comm->config.nChannelsPerNetPeer = NCCL_CONFIG_UNDEF_INT; + comm->nNodes = 1; + EXPECT_EQ(2, rcclP2pTaskParts(comm.get(), ncclFuncScatter)) << "single node"; +} + +TEST_F(EnqueueMicrotest, P2pPolicyChannels_ScatterBelowPerPeerReturnsItsParts) { + // Returns before the arch policy table, so comm->topo is never read. + g_rcclUseAinicValue = true; + auto comm = P2pPartsComm(/*nNodes=*/4, /*nRanks=*/32, /*p2pnChannels=*/32, /*perPeer=*/2); + ncclTaskP2p task{}; + task.collAPI = ncclFuncScatter; + EXPECT_EQ(1, rcclP2pPolicyChannels(comm.get(), &task)); + EXPECT_EQ(-1, rcclP2pPolicyChannels(comm.get(), nullptr)); +} + // =========================================================================== // calcP2pChannelCount (enqueue.cc:1550) // diff --git a/projects/rccl/test/host/fakes/enqueue_fakes.cc b/projects/rccl/test/host/fakes/enqueue_fakes.cc index eaae182389ee..1071fa886ca6 100644 --- a/projects/rccl/test/host/fakes/enqueue_fakes.cc +++ b/projects/rccl/test/host/fakes/enqueue_fakes.cc @@ -17,7 +17,7 @@ #include "signature-drift.h" -#include "channel.h" // rcclEffectiveP2pBatchEnable +#include "channel.h" // rcclEffectiveP2pBatchEnable, rcclP2pTaskParts #include "enqueue_fakes.h" #include "nccl_fakes.h" // g_loadParam, for the NCCL_PARAM body this file stands in for #include "sym_kernels_fakes.h" // g_symkAvailable and the sym_kernels.cc seams below's canonical home @@ -31,6 +31,7 @@ ASSERT_HOOK_MATCHES_PROD(g_addProxyOpIfNeeded, ncclAddProxyOpIfNeeded); ASSERT_HOOK_MATCHES_PROD(g_getCollNetSupport, ncclGetCollNetSupport); ASSERT_HOOK_MATCHES_PROD(g_getRegBuff, ncclGetRegBuff); ASSERT_HOOK_MATCHES_PROD(g_rcclEffectiveP2pBatchEnable, rcclEffectiveP2pBatchEnable); +ASSERT_HOOK_MATCHES_PROD(g_rcclP2pTaskParts, rcclP2pTaskParts); #undef ASSERT_HOOK_MATCHES_PROD // Generous default: a deny-everything default would make even a single small task look over budget. @@ -75,6 +76,10 @@ std::function g_rcclEffectiveP2pBatchEnable = DefaultEffectiveP2pBatchEnable; +// enqueue.cc's answer for every task outside multi-node AINIC Scatter. +static int DefaultP2pTaskParts(struct ncclComm* comm, ncclFunc_t) { return comm->p2pnChannelsPerPeer; } +std::function g_rcclP2pTaskParts = DefaultP2pTaskParts; + // enqueue.cc:437-441 splits the default on ROCm version; mirrored rather than pinned to one arm. #if ROCM_VERSION >= 71200 int64_t ncclParamGraphRegister() { return g_loadParam("GRAPH_REGISTER", 1); } @@ -105,11 +110,15 @@ void ResetEnqueueFakes() { g_getCollNetSupport = DefaultGetCollNetSupport; g_getRegBuff = DefaultGetRegBuff; g_rcclEffectiveP2pBatchEnable = DefaultEffectiveP2pBatchEnable; + g_rcclP2pTaskParts = DefaultP2pTaskParts; } int rcclEffectiveP2pBatchEnable(struct ncclComm* comm) { return g_rcclEffectiveP2pBatchEnable(comm); } +int rcclP2pTaskParts(struct ncclComm* comm, ncclFunc_t collAPI) { + return g_rcclP2pTaskParts(comm, collAPI); +} // src/enqueue/enqueue.cc ncclResult_t ncclGetAlgoInfo(struct ncclComm* comm, struct ncclTaskColl* task, int collNetSupport, int nvlsSupport, diff --git a/projects/rccl/test/host/fakes/enqueue_fakes.h b/projects/rccl/test/host/fakes/enqueue_fakes.h index 4acd428b1e81..88b6ccca3c8d 100644 --- a/projects/rccl/test/host/fakes/enqueue_fakes.h +++ b/projects/rccl/test/host/fakes/enqueue_fakes.h @@ -52,6 +52,9 @@ extern std::function // enqueue.cc's rcclEffectiveP2pBatchEnable: default 0 is batching off, the single-node answer. extern std::function g_rcclEffectiveP2pBatchEnable; +// enqueue.cc's rcclP2pTaskParts: default p2pnChannelsPerPeer; install the multi-node AINIC Scatter limit per test. +extern std::function g_rcclP2pTaskParts; + void ResetEnqueueFakes(); #endif // RCCL_TEST_HOST_FAKES_ENQUEUE_FAKES_H_ diff --git a/projects/rccl/test/host/task-posttuning-test.cc b/projects/rccl/test/host/task-posttuning-test.cc index 25910c02bbd5..1b88f66e725b 100644 --- a/projects/rccl/test/host/task-posttuning-test.cc +++ b/projects/rccl/test/host/task-posttuning-test.cc @@ -1303,13 +1303,13 @@ TEST_F(TaskPostTuningMicrotest, P2pRecordPreconnect_PeerOutsideTheCommunicator_R for (int peer : {-1, kRanks}) { bool needPreconnect = false; - EXPECT_EQ(ncclInvalidArgument, postTuneP2pRecordPreconnect(p2p.comm(), peer, kIsSend, &needPreconnect)) + EXPECT_EQ(ncclInvalidArgument, postTuneP2pRecordPreconnect(p2p.comm(), peer, kIsSend, ncclFuncSendRecv, &needPreconnect)) << "peer " << peer; EXPECT_FALSE(needPreconnect) << "peer " << peer; needPreconnect = true; - EXPECT_EQ(ncclInvalidArgument, postTuneP2pRecordPreconnect(p2p.comm(), peer, kIsSend, &needPreconnect)) + EXPECT_EQ(ncclInvalidArgument, postTuneP2pRecordPreconnect(p2p.comm(), peer, kIsSend, ncclFuncSendRecv, &needPreconnect)) << "peer " << peer; EXPECT_TRUE(needPreconnect) << "peer " << peer; @@ -1320,14 +1320,14 @@ TEST_F(TaskPostTuningMicrotest, P2pRecordPreconnect_PeerIsThisRank_SucceedsWitho TaskPostTuning_P2pScene p2p; bool needPreconnect = false; - EXPECT_EQ(ncclSuccess, postTuneP2pRecordPreconnect(p2p.comm(), kP2pRank, kIsSend, &needPreconnect)); + EXPECT_EQ(ncclSuccess, postTuneP2pRecordPreconnect(p2p.comm(), kP2pRank, kIsSend, ncclFuncSendRecv, &needPreconnect)); EXPECT_FALSE(needPreconnect); EXPECT_FALSE(p2p.comm()->planner.peers[kP2pRank].sendSeen); EXPECT_TRUE(TaskPostTuning_MaskHasExactly(p2p.connectSend(kP2pRank), {})); needPreconnect = true; - EXPECT_EQ(ncclSuccess, postTuneP2pRecordPreconnect(p2p.comm(), kP2pRank, kIsSend, &needPreconnect)); + EXPECT_EQ(ncclSuccess, postTuneP2pRecordPreconnect(p2p.comm(), kP2pRank, kIsSend, ncclFuncSendRecv, &needPreconnect)); EXPECT_TRUE(needPreconnect); } @@ -1337,7 +1337,7 @@ TEST_F(TaskPostTuningMicrotest, P2pRecordPreconnect_FreshSendPeer_ClaimsItsSendC const std::vector sendChannels = p2p.Channels(kIsSend); bool needPreconnect = false; - ASSERT_EQ(ncclSuccess, postTuneP2pRecordPreconnect(p2p.comm(), kP2pPeer, kIsSend, &needPreconnect)); + ASSERT_EQ(ncclSuccess, postTuneP2pRecordPreconnect(p2p.comm(), kP2pPeer, kIsSend, ncclFuncSendRecv, &needPreconnect)); EXPECT_TRUE(needPreconnect); EXPECT_TRUE(p2p.comm()->planner.peers[kP2pPeer].sendSeen); @@ -1359,7 +1359,7 @@ TEST_F(TaskPostTuningMicrotest, P2pRecordPreconnect_FreshRecvPeer_ClaimsItsRecvC const std::vector recvChannels = p2p.Channels(kIsRecv); bool needPreconnect = false; - ASSERT_EQ(ncclSuccess, postTuneP2pRecordPreconnect(p2p.comm(), kP2pPeer, kIsRecv, &needPreconnect)); + ASSERT_EQ(ncclSuccess, postTuneP2pRecordPreconnect(p2p.comm(), kP2pPeer, kIsRecv, ncclFuncSendRecv, &needPreconnect)); EXPECT_TRUE(needPreconnect); EXPECT_TRUE(p2p.comm()->planner.peers[kP2pPeer].recvSeen); @@ -1376,6 +1376,37 @@ TEST_F(TaskPostTuningMicrotest, P2pRecordPreconnect_FreshRecvPeer_ClaimsItsRecvC EXPECT_TRUE(TaskPostTuning_MaskHasExactly(p2p.connectRecv(kP2pRank), {})); } +// Scatter with fewer parts than p2pnChannelsPerPeer claims only its own channels and leaves the side +// unseen, so a later task to the same peer that uses every part still claims the rest. +TEST_F(TaskPostTuningMicrotest, P2pRecordPreconnect_ScatterWithFewerParts_ClaimsItsPartsAndLeavesTheRestForLater) { + for (bool isSendNotRecv : {kIsSend, kIsRecv}) { + TaskPostTuning_P2pScene p2p; + ScopedHook scatterParts(g_rcclP2pTaskParts, [](struct ncclComm* comm, ncclFunc_t collAPI) { + return collAPI == ncclFuncScatter ? 1 : comm->p2pnChannelsPerPeer; + }); + const std::vector channels = p2p.Channels(isSendNotRecv); + ASSERT_EQ(2u, channels.size()); + bool needPreconnect = false; + + ASSERT_EQ(ncclSuccess, + postTuneP2pRecordPreconnect(p2p.comm(), kP2pPeer, isSendNotRecv, ncclFuncScatter, &needPreconnect)); + + EXPECT_TRUE(needPreconnect); + EXPECT_FALSE(p2p.Seen(isSendNotRecv, kP2pPeer)); + EXPECT_EQ(kConnected, p2p.Conn(isSendNotRecv, channels[0], kP2pPeer)->hasSeen); + EXPECT_EQ(kNotConnected, p2p.Conn(isSendNotRecv, channels[1], kP2pPeer)->hasSeen); + EXPECT_TRUE(TaskPostTuning_MaskHasExactly(p2p.connect(isSendNotRecv, kP2pPeer), {channels[0]})); + + // The full-part task adds only the channel the Scatter left out. + ASSERT_EQ(ncclSuccess, + postTuneP2pRecordPreconnect(p2p.comm(), kP2pPeer, isSendNotRecv, ncclFuncSendRecv, &needPreconnect)); + + EXPECT_TRUE(p2p.Seen(isSendNotRecv, kP2pPeer)); + EXPECT_EQ(kConnected, p2p.Conn(isSendNotRecv, channels[1], kP2pPeer)->hasSeen); + EXPECT_TRUE(TaskPostTuning_MaskHasExactly(p2p.connect(isSendNotRecv, kP2pPeer), channels)); + } +} + // Production duplicates the loop body per side, so each of these runs on the recv copy as well as the send one. TEST_F(TaskPostTuningMicrotest, P2pRecordPreconnect_PeerAlreadyRecordedOnThatSide_ClaimsNoChannel) { for (bool isSendNotRecv : {kIsSend, kIsRecv}) { @@ -1383,7 +1414,7 @@ TEST_F(TaskPostTuningMicrotest, P2pRecordPreconnect_PeerAlreadyRecordedOnThatSid p2p.Seen(isSendNotRecv, kP2pPeer) = true; bool needPreconnect = false; - EXPECT_EQ(ncclSuccess, postTuneP2pRecordPreconnect(p2p.comm(), kP2pPeer, isSendNotRecv, &needPreconnect)) + EXPECT_EQ(ncclSuccess, postTuneP2pRecordPreconnect(p2p.comm(), kP2pPeer, isSendNotRecv, ncclFuncSendRecv, &needPreconnect)) << "isSendNotRecv " << isSendNotRecv; EXPECT_FALSE(needPreconnect) << "isSendNotRecv " << isSendNotRecv; @@ -1394,7 +1425,7 @@ TEST_F(TaskPostTuningMicrotest, P2pRecordPreconnect_PeerAlreadyRecordedOnThatSid } needPreconnect = true; - EXPECT_EQ(ncclSuccess, postTuneP2pRecordPreconnect(p2p.comm(), kP2pPeer, isSendNotRecv, &needPreconnect)) + EXPECT_EQ(ncclSuccess, postTuneP2pRecordPreconnect(p2p.comm(), kP2pPeer, isSendNotRecv, ncclFuncSendRecv, &needPreconnect)) << "isSendNotRecv " << isSendNotRecv; EXPECT_TRUE(needPreconnect) << "isSendNotRecv " << isSendNotRecv; @@ -1406,7 +1437,7 @@ TEST_F(TaskPostTuningMicrotest, P2pRecordPreconnect_PeerRecordedOnTheOtherSideOn p2p.comm()->planner.peers[kP2pPeer].recvSeen = true; bool needPreconnect = false; - ASSERT_EQ(ncclSuccess, postTuneP2pRecordPreconnect(p2p.comm(), kP2pPeer, kIsSend, &needPreconnect)); + ASSERT_EQ(ncclSuccess, postTuneP2pRecordPreconnect(p2p.comm(), kP2pPeer, kIsSend, ncclFuncSendRecv, &needPreconnect)); EXPECT_TRUE(needPreconnect); EXPECT_TRUE(TaskPostTuning_MaskHasExactly(p2p.connectSend(kP2pPeer), @@ -1420,7 +1451,7 @@ TEST_F(TaskPostTuningMicrotest, P2pRecordPreconnect_SomeChannelsAlreadyConnected p2p.Conn(isSendNotRecv, channels.front(), kP2pPeer)->hasSeen = kConnected; bool needPreconnect = false; - ASSERT_EQ(ncclSuccess, postTuneP2pRecordPreconnect(p2p.comm(), kP2pPeer, isSendNotRecv, &needPreconnect)) + ASSERT_EQ(ncclSuccess, postTuneP2pRecordPreconnect(p2p.comm(), kP2pPeer, isSendNotRecv, ncclFuncSendRecv, &needPreconnect)) << "isSendNotRecv " << isSendNotRecv; EXPECT_TRUE(needPreconnect) << "isSendNotRecv " << isSendNotRecv; @@ -1439,7 +1470,7 @@ TEST_F(TaskPostTuningMicrotest, P2pRecordPreconnect_EveryChannelAlreadyConnected } bool needPreconnect = false; - ASSERT_EQ(ncclSuccess, postTuneP2pRecordPreconnect(p2p.comm(), kP2pPeer, isSendNotRecv, &needPreconnect)) + ASSERT_EQ(ncclSuccess, postTuneP2pRecordPreconnect(p2p.comm(), kP2pPeer, isSendNotRecv, ncclFuncSendRecv, &needPreconnect)) << "isSendNotRecv " << isSendNotRecv; EXPECT_FALSE(needPreconnect) << "isSendNotRecv " << isSendNotRecv; @@ -1457,7 +1488,7 @@ TEST_F(TaskPostTuningMicrotest, P2pRecordPreconnect_CallerAlreadyNeedsSetup_Keep } bool needPreconnect = true; - ASSERT_EQ(ncclSuccess, postTuneP2pRecordPreconnect(p2p.comm(), kP2pPeer, isSendNotRecv, &needPreconnect)) + ASSERT_EQ(ncclSuccess, postTuneP2pRecordPreconnect(p2p.comm(), kP2pPeer, isSendNotRecv, ncclFuncSendRecv, &needPreconnect)) << "isSendNotRecv " << isSendNotRecv; EXPECT_TRUE(needPreconnect) << "isSendNotRecv " << isSendNotRecv; @@ -1470,7 +1501,7 @@ TEST_F(TaskPostTuningMicrotest, P2pRecordPreconnect_ChannelsBeyondTheFirstMaskWo const std::vector channels = p2p.Channels(isSendNotRecv); bool needPreconnect = false; - ASSERT_EQ(ncclSuccess, postTuneP2pRecordPreconnect(p2p.comm(), kP2pPeer, isSendNotRecv, &needPreconnect)) + ASSERT_EQ(ncclSuccess, postTuneP2pRecordPreconnect(p2p.comm(), kP2pPeer, isSendNotRecv, ncclFuncSendRecv, &needPreconnect)) << "isSendNotRecv " << isSendNotRecv; EXPECT_LE(CHANNELS_PER_MASK_WORD, channels.front()) << "isSendNotRecv " << isSendNotRecv; @@ -1486,7 +1517,7 @@ TEST_F(TaskPostTuningMicrotest, P2pRecordPreconnect_ChannelRangeStraddlingAMaskW const std::vector channels = p2p.Channels(isSendNotRecv); bool needPreconnect = false; - ASSERT_EQ(ncclSuccess, postTuneP2pRecordPreconnect(p2p.comm(), kP2pPeer, isSendNotRecv, &needPreconnect)) + ASSERT_EQ(ncclSuccess, postTuneP2pRecordPreconnect(p2p.comm(), kP2pPeer, isSendNotRecv, ncclFuncSendRecv, &needPreconnect)) << "isSendNotRecv " << isSendNotRecv; EXPECT_NE(0, channels.front() % CHANNELS_PER_MASK_WORD) << "isSendNotRecv " << isSendNotRecv;