From 5dc0670e56fbd0ec2290dbfb5af5d1428d49a39e Mon Sep 17 00:00:00 2001
From: David Addison <daddison@nvidia.com>
Date: Mon, 30 Mar 2026 11:39:31 -0700
Subject: [PATCH] Add maxP2pPeers comm config for sendrecv

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>
---
 src/common.cu   |  2 ++
 src/common.h    |  4 ++++
 src/sendrecv.cu | 16 +++++++++++++++-
 3 files changed, 21 insertions(+), 1 deletion(-)

diff --git a/src/common.cu b/src/common.cu
index 9dff7af..b14c3d5 100644
--- a/src/common.cu
+++ b/src/common.cu
@@ -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());
diff --git a/src/common.h b/src/common.h
index a68ecaa..a342057 100644
--- a/src/common.h
+++ b/src/common.h
@@ -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);
diff --git a/src/sendrecv.cu b/src/sendrecv.cu
index 2f09667..66e8af0 100644
--- a/src/sendrecv.cu
+++ b/src/sendrecv.cu
@@ -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