kanshan/nccl-tests
1
0
Fork 0
mirror of https://github.com/NVIDIA/nccl-tests.git synced 2026-04-23 16:08:20 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

Add maxP2pPeers comm config for sendrecv

Browse Source 
Add optional testEngine.initCommConfig, invoked from initComms
after the shared ncclConfig_t setup.

sendrecv registers SendRecvInitCommConfig to set maxP2pPeers=2

Signed-off-by: David Addison <daddison@nvidia.com>
This commit is contained in:
David Addison 2026-03-30 11:39:31 -07:00
parent af1dcac92a
commit 5dc0670e56
3 changed files with 21 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
src/common.cu
View File

@ -195,6 +195,8 @@ testResult_t initComms(ncclComm_t* comms, int nComms, int firstRank, int nRanks,
config.nvlinkCentricSched = 1;
#endif
#endif
if (ncclTestEngine.initCommConfig)
ncclTestEngine.initCommConfig(&config);
#endif
NCCLCHECK(ncclGroupStart());

4
src/common.h
View File

@ -112,6 +112,10 @@ struct testEngine {
void (*getBuffSize)(size_t *sendcount, size_t *recvcount, size_t count, int nranks);
testResult_t (*runTest)(struct threadArgs* args, int root, ncclDataType_t type,
const char* typeName, ncclRedOp_t op, const char* opName);
#if NCCL_VERSION_CODE >= NCCL_VERSION(2,14,0)
/* Optional; called from initComms after common fields are set on ncclConfig_t. */
void (*initCommConfig)(ncclConfig_t* config);
#endif
#if NCCL_VERSION_CODE >= NCCL_VERSION(2,29,0)
testResult_t (*getDevCommRequirements)(int deviceImpl, ncclDevCommRequirements* reqs, ncclCommProperties_t* commProperties);

16
src/sendrecv.cu
View File

@ -64,6 +64,17 @@ testResult_t SendRecvRunColl(void* sendbuff, size_t sendoffset, void* recvbuff,
return testSuccess;
}
#if NCCL_VERSION_CODE >= NCCL_VERSION(2,14,0)
static void SendRecvInitCommConfig(ncclConfig_t* config) {
#if NCCL_VERSION_CODE >= NCCL_VERSION(2,30,0)
// Ring sendrecv uses two distinct peers per rank (send to rank+1, recv from rank-1).
config->maxP2pPeers = 2;
#else
(void)config;
#endif
}
#endif
struct testColl sendRecvTest = {
"SendRecv",
SendRecvGetCollByteCount,
@ -114,7 +125,10 @@ testResult_t SendRecvRunTest(struct threadArgs* args, int root, ncclDataType_t t
struct testEngine sendRecvEngine = {
.getBuffSize = SendRecvGetBuffSize,
.runTest = SendRecvRunTest
.runTest = SendRecvRunTest,
#if NCCL_VERSION_CODE >= NCCL_VERSION(2,14,0)
.initCommConfig = SendRecvInitCommConfig,
#endif
};
#pragma weak ncclTestEngine=sendRecvEngine