From 492834789a7c6daa5e2ff08686bd04d5e8c4a859 Mon Sep 17 00:00:00 2001
From: Devendar Bureddy <devendar@nvidia.com>
Date: Thu, 26 Mar 2026 00:30:11 +0000
Subject: [PATCH] Support collect nic-fusion

Refactor device name processing in SHARP plugin to support merged NICs and direct NIC identification. Update device property retrieval to use the new v11 interface. Add makeVDevice functions for v10 and v11 compatibility.
---
 src/sharp_plugin.c | 69 ++++++++++++++++++++++++++++++----------------
 1 file changed, 45 insertions(+), 24 deletions(-)

diff --git a/src/sharp_plugin.c b/src/sharp_plugin.c
index c86b7df..807e5e3 100644
--- a/src/sharp_plugin.c
+++ b/src/sharp_plugin.c
@@ -265,8 +265,7 @@ ncclResult_t ncclSharpInit(void **ctx, uint64_t commId, ncclDebugLogger_t logFun
 }
 
 ncclResult_t ncclSharpDevices(int* ndev) {
-  *ndev = ncclNSharpDevs;
-  return ncclSuccess;
+  return ncclNetPlugin_v11.devices(ndev);
 }
 
 ncclResult_t ncclSharpGetProperties_v11(int dev, ncclNetProperties_v11_t* props) {
@@ -318,24 +317,39 @@ ncclResult_t ncclSharpListen(void *ctx, int dev, void* opaqueHandle, void** list
   return status;
 }
 
-/* for data direct nic, the device name is ends with suffix '_dma`.
- * remove this suffix before passing name to libsharp */
-static void sharp_get_device_name(const char *input, char *output, size_t output_size) {
+/* Build libsharp device list: split merged NICs on '+', strip optional
+ * '_dma' suffix per token (data-direct), join with ',', append ':1' per device. */
+static void ncclSharpProcessDeviceName(const char *input, char *output, size_t output_size, int *isDirectNic) {
   const char *suffix = "_dma";
-  size_t input_len = strlen(input);
-  size_t suffix_len = strlen(suffix);
-
-  if (input_len >= suffix_len && strcmp(input + input_len - suffix_len, suffix) == 0) {
-    size_t new_len = input_len - suffix_len;
-    if (new_len >= output_size) {
-      new_len = output_size - 1;
+  const size_t suffix_len = strlen(suffix);
+  *isDirectNic = 0;
+  if (output_size == 0) return;
+  output[0] = '\0';
+
+  char buf[MAX_MERGED_DEV_NAME];
+  strncpy(buf, input, sizeof(buf) - 1);
+  buf[sizeof(buf) - 1] = '\0';
+
+  size_t o = 0;
+  char *saveptr;
+  for (char *tok = strtok_r(buf, "+", &saveptr); tok; tok = strtok_r(NULL, "+", &saveptr)) {
+    size_t len = strlen(tok);
+    if (len >= suffix_len && memcmp(tok + len - suffix_len, suffix, suffix_len) == 0) {
+      len -= suffix_len;
+      *isDirectNic = 1;
     }
-    memcpy(output, input, new_len);
-    output[new_len] = '\0';
-  } else {
-    strncpy(output, input, output_size - 1);
-    output[output_size - 1] = '\0';
+    if (len == 0) continue;
+
+    size_t need = (o > 0 ? 1 : 0) + len + 2 + 1;
+    if (o + need > output_size) break;
+
+    if (o > 0) output[o++] = ',';
+    memcpy(output + o, tok, len);
+    o += len;
+    memcpy(output + o, ":1", 2);
+    o += 2;
   }
+  output[o] = '\0';
 }
 
 ncclResult_t ncclSharpConnect(void* handles[], int nranks, int rank, void* listenComm, void** collComm) {
@@ -410,11 +424,11 @@ ncclResult_t ncclSharpConnect(void* handles[], int nranks, int rank, void* liste
   init_spec.config = sharp_coll_default_config;
   init_spec.config.user_progress_num_polls = 10000000;
 
-  char devName[MAXNAMESIZE];
-  ncclNetProperties_v6_t prop;
-  ncclSharpGetProperties_v6(lComm->dev, &prop);
-  sharp_get_device_name(prop.name, devName, 64);
-  snprintf(devName + strlen(devName), MAXNAMESIZE - strlen(devName), ":%d", prop.port);
+  char devName[MAX_MERGED_DEV_NAME];
+  int isDirectNic = 0;
+  ncclNetProperties_v11_t prop;
+  ncclSharpGetProperties_v11(lComm->dev, &prop);
+  ncclSharpProcessDeviceName(prop.name, devName, sizeof(devName), &isDirectNic);
   init_spec.config.ib_dev_list = devName;
 
   int ret = sharp_coll_init(&init_spec, &cComm->sharpCollContext);
@@ -850,7 +864,14 @@ ncclResult_t ncclSharpInit_v10(ncclDebugLogger_t logFunction) {
 ncclResult_t ncclSharpListen_v10(int dev, void* opaqueHandle, void** listenComm) {
   return ncclSharpListen(NULL, dev, opaqueHandle, listenComm);
 }
- 
+
+ncclResult_t ncclSharpMakeVDevice(int* d, ncclNetVDeviceProps_v11_t* props) {
+  return ncclNetPlugin_v11.makeVDevice(d, props);
+}
+
+ncclResult_t ncclSharpMakeVDevice_v10(int* d, ncclNetVDeviceProps_v10_t* props) {
+  return ncclNetPlugin_v10.makeVDevice(d, props);
+}
 
 ncclCollNet_v11_t ncclCollNetPlugin_v11 = {
   "SHARP",
@@ -870,7 +891,7 @@ ncclCollNet_v11_t ncclCollNetPlugin_v11 = {
   ncclSharpTest,
   ncclSharpCloseColl,
   ncclSharpCloseListen,
-  NULL,
+  ncclSharpMakeVDevice,
   ncclSharpFinalize
 };