GIN perf: Rename 'ginRsm_t' and 'ginOp_t'

Rename 'ginRsm_t' to 'ncclGinRsm_t' and 'ginOp_t' to 'ncclGinOp_t',
aslo rename the emum in it to comply the coding style.

Signed-off-by: Bojun Zhu <[email protected]>

src/device_api/gin/latency/ginGetLatency_ping.cu

@@ -39,9 +39,9 @@ static void ginGetLatencyPingLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle
 
 void ncclTestGinGetLatencyPingLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
                                      const ginArgs_t* args, size_t numElems, int iters) {
-  if (args->ginRsm == GIN_RSM_THREAD) {
-    ginGetLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_THREAD>(dcomm, hBuf, stream, args, numElems, iters);
-  } else if (args->ginRsm == GIN_RSM_CTA) {
+  if (args->ginRsm == ncclGinRsmThread) {
+    ginGetLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_THREAD>(dcomm, devBufHandle, hostBufWindow, stream, args, numElems, iters);
+  } else if (args->ginRsm == ncclGinRsmCta) {
     ginGetLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, hBuf, stream, args, numElems, iters);
   } else {
     ginGetLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, hBuf, stream, args, numElems, iters);

src/device_api/gin/latency/ginGetLatency_ping_main.cu

@@ -8,10 +8,10 @@ int main(int argc, char** argv) {
   ncclTestGinParseArgs(argc, argv, &args);
 
   ginTestCaps_t caps = {};
-  caps.defaultRsm = GIN_RSM_GPU;
+  caps.defaultRsm = ncclGinRsmGpu;
   caps.allowThreadRsm = true;
   caps.opMask = 0;
-  caps.defaultOp = GIN_OP_UNSET;
+  caps.defaultOp = ncclGinOpUnset;
   caps.isSignalOp = false;
   caps.allowAggregateRequests = false;
   caps.allowBidirectional = false;

src/device_api/gin/latency/ginPutCounterLatency_ping.cu

@@ -50,9 +50,9 @@ static void ginPutCounterLatencyPingLaunchRsm(ncclDevComm dcomm, ncclDevResource
 
 void ncclTestGinPutCounterLatencyPingLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
                                             const ginArgs_t* args, size_t numElems, int iters) {
-  if (args->ginRsm == GIN_RSM_THREAD) {
+  if (args->ginRsm == ncclGinRsmThread) {
     ginPutCounterLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_THREAD>(dcomm, hBuf, stream, args, numElems, iters);
-  } else if (args->ginRsm == GIN_RSM_CTA) {
+  } else if (args->ginRsm == ncclGinRsmCta) {
     ginPutCounterLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, hBuf, stream, args, numElems, iters);
   } else {
     ginPutCounterLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, hBuf, stream, args, numElems, iters);

src/device_api/gin/latency/ginPutLatency_ping.cu

@@ -39,9 +39,9 @@ static void ginPutLatencyPingLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle
 
 void ncclTestGinPutLatencyPingLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
                                      const ginArgs_t* args, size_t numElems, int iters) {
-  if (args->ginRsm == GIN_RSM_THREAD) {
+  if (args->ginRsm == ncclGinRsmThread) {
     ginPutLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_THREAD>(dcomm, hBuf, stream, args, numElems, iters);
-  } else if (args->ginRsm == GIN_RSM_CTA) {
+  } else if (args->ginRsm == ncclGinRsmCta) {
     ginPutLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, hBuf, stream, args, numElems, iters);
   } else {
     ginPutLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, hBuf, stream, args, numElems, iters);

src/device_api/gin/latency/ginPutLatency_pingPong.cu

@@ -62,9 +62,9 @@ static void ginPutLatencyPingPongLaunchRsm(ncclDevComm dcomm, ncclDevResourceHan
 
 void ncclTestGinPutLatencyPingPongLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
                                          const ginArgs_t* args, size_t numElems, int iters) {
-  if (args->ginRsm == GIN_RSM_THREAD) {
+  if (args->ginRsm == ncclGinRsmThread) {
     ginPutLatencyPingPongLaunchRsm<NCCL_GIN_RESOURCE_SHARING_THREAD>(dcomm, hBuf, stream, args, numElems, iters);
-  } else if (args->ginRsm == GIN_RSM_CTA) {
+  } else if (args->ginRsm == ncclGinRsmCta) {
     ginPutLatencyPingPongLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, hBuf, stream, args, numElems, iters);
   } else {
     ginPutLatencyPingPongLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, hBuf, stream, args, numElems, iters);

src/device_api/gin/latency/ginPutLatency_pingPong_main.cu

@@ -11,10 +11,10 @@ int main(int argc, char** argv) {
   ncclTestGinParseArgs(argc, argv, &args);
 
   ginTestCaps_t caps = {};
-  caps.defaultRsm = GIN_RSM_GPU;
+  caps.defaultRsm = ncclGinRsmGpu;
   caps.allowThreadRsm = true;
-  caps.opMask = GIN_OP_BIT(GIN_OP_PUT) | GIN_OP_BIT(GIN_OP_PUT_SIGNAL);
-  caps.defaultOp = GIN_OP_PUT;
+  caps.opMask = GIN_OP_BIT(ncclGinOpPut) | GIN_OP_BIT(ncclGinOpPutSignal);
+  caps.defaultOp = ncclGinOpPut;
   caps.isSignalOp = false;
   caps.allowAggregateRequests = false;
   caps.allowBidirectional = false;
@@ -24,7 +24,7 @@ int main(int argc, char** argv) {
   ginBenchmark_t bench = {};
   bench.type = GIN_BENCHMARK_TYPE_PING_PONG;
   switch (args.ginOp) {
-  case GIN_OP_PUT_SIGNAL:
+  case ncclGinOpPutSignal:
     bench.run = ncclTestGinPutSignalLatencyPingPongLaunch;
     bench.name = "GIN PUT_SIGNAL PING-PONG LATENCY";
     break;

src/device_api/gin/latency/ginPutLatency_ping_main.cu

@@ -8,10 +8,10 @@ int main(int argc, char** argv) {
   ncclTestGinParseArgs(argc, argv, &args);
 
   ginTestCaps_t caps = {};
-  caps.defaultRsm = GIN_RSM_GPU;
+  caps.defaultRsm = ncclGinRsmGpu;
   caps.allowThreadRsm = true;
-  caps.opMask = GIN_OP_BIT(GIN_OP_PUT) | GIN_OP_BIT(GIN_OP_PUT_SIGNAL) | GIN_OP_BIT(GIN_OP_PUT_COUNTER);
-  caps.defaultOp = GIN_OP_PUT;
+  caps.opMask = GIN_OP_BIT(ncclGinOpPut) | GIN_OP_BIT(ncclGinOpPutSignal) | GIN_OP_BIT(ncclGinOpPutCount);
+  caps.defaultOp = ncclGinOpPut;
   caps.isSignalOp = false;
   caps.allowAggregateRequests = false;
   caps.allowBidirectional = false;
@@ -21,11 +21,11 @@ int main(int argc, char** argv) {
   ginBenchmark_t bench = {};
   bench.type = GIN_BENCHMARK_TYPE_PING;
   switch (args.ginOp) {
-  case GIN_OP_PUT_SIGNAL:
+  case ncclGinOpPutSignal:
     bench.run = ncclTestGinPutSignalLatencyPingLaunch;
     bench.name = "GIN PUT_SIGNAL PING LATENCY";
     break;
-  case GIN_OP_PUT_COUNTER:
+  case ncclGinOpPutCount:
     bench.run = ncclTestGinPutCounterLatencyPingLaunch;
     bench.name = "GIN PUT_COUNTER PING LATENCY";
     break;

src/device_api/gin/latency/ginPutSignalLatency_ping.cu

@@ -51,9 +51,9 @@ static void ginPutSignalLatencyPingLaunchRsm(ncclDevComm dcomm, ncclDevResourceH
 
 void ncclTestGinPutSignalLatencyPingLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
                                            const ginArgs_t* args, size_t numElems, int iters) {
-  if (args->ginRsm == GIN_RSM_THREAD) {
+  if (args->ginRsm == ncclGinRsmThread) {
     ginPutSignalLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_THREAD>(dcomm, hBuf, stream, args, numElems, iters);
-  } else if (args->ginRsm == GIN_RSM_CTA) {
+  } else if (args->ginRsm == ncclGinRsmCta) {
     ginPutSignalLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, hBuf, stream, args, numElems, iters);
   } else {
     ginPutSignalLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, hBuf, stream, args, numElems, iters);

src/device_api/gin/latency/ginPutSignalLatency_pingPong.cu

@@ -72,9 +72,9 @@ static void ginPutSignalLatencyPingPongLaunchRsm(ncclDevComm dcomm, ncclDevResou
 
 void ncclTestGinPutSignalLatencyPingPongLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
                                                const ginArgs_t* args, size_t numElems, int iters) {
-  if (args->ginRsm == GIN_RSM_THREAD) {
+  if (args->ginRsm == ncclGinRsmThread) {
     ginPutSignalLatencyPingPongLaunchRsm<NCCL_GIN_RESOURCE_SHARING_THREAD>(dcomm, hBuf, stream, args, numElems, iters);
-  } else if (args->ginRsm == GIN_RSM_CTA) {
+  } else if (args->ginRsm == ncclGinRsmCta) {
     ginPutSignalLatencyPingPongLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, hBuf, stream, args, numElems, iters);
   } else {
     ginPutSignalLatencyPingPongLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, hBuf, stream, args, numElems, iters);

src/device_api/gin/latency/ginSignalLatency_ping.cu

@@ -49,9 +49,9 @@ static void ginSignalLatencyPingLaunchRsm(ncclDevComm dcomm, ncclDevResourceHand
 
 void ncclTestGinSignalLatencyPingLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
                                         const ginArgs_t* args, size_t numElems, int iters) {
-  if (args->ginRsm == GIN_RSM_THREAD) {
+  if (args->ginRsm == ncclGinRsmThread) {
     ginSignalLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_THREAD>(dcomm, hBuf, stream, args, iters);
-  } else if (args->ginRsm == GIN_RSM_CTA) {
+  } else if (args->ginRsm == ncclGinRsmCta) {
     ginSignalLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, hBuf, stream, args, iters);
   } else {
     ginSignalLatencyPingLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, hBuf, stream, args, iters);

src/device_api/gin/latency/ginSignalLatency_pingPong.cu

@@ -70,9 +70,9 @@ static void ginSignalLatencyPingPongLaunchRsm(ncclDevComm dcomm, ncclDevResource
 
 void ncclTestGinSignalLatencyPingPongLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
                                             const ginArgs_t* args, size_t numElems, int iters) {
-  if (args->ginRsm == GIN_RSM_THREAD) {
+  if (args->ginRsm == ncclGinRsmThread) {
     ginSignalLatencyPingPongLaunchRsm<NCCL_GIN_RESOURCE_SHARING_THREAD>(dcomm, hBuf, stream, args, iters);
-  } else if (args->ginRsm == GIN_RSM_CTA) {
+  } else if (args->ginRsm == ncclGinRsmCta) {
     ginSignalLatencyPingPongLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, hBuf, stream, args, iters);
   } else {
     ginSignalLatencyPingPongLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, hBuf, stream, args, iters);

src/device_api/gin/latency/ginSignalLatency_pingPong_main.cu

@@ -8,10 +8,10 @@ int main(int argc, char** argv) {
   ncclTestGinParseArgs(argc, argv, &args);
 
   ginTestCaps_t caps = {};
-  caps.defaultRsm = GIN_RSM_GPU;
+  caps.defaultRsm = ncclGinRsmGpu;
   caps.allowThreadRsm = true;
   caps.opMask = 0;
-  caps.defaultOp = GIN_OP_UNSET;
+  caps.defaultOp = ncclGinOpUnset;
   caps.isSignalOp = true;
   caps.allowAggregateRequests = false;
   caps.allowBidirectional = false;

src/device_api/gin/latency/ginSignalLatency_ping_main.cu

@@ -8,10 +8,10 @@ int main(int argc, char** argv) {
   ncclTestGinParseArgs(argc, argv, &args);
 
   ginTestCaps_t caps = {};
-  caps.defaultRsm = GIN_RSM_GPU;
+  caps.defaultRsm = ncclGinRsmGpu;
   caps.allowThreadRsm = true;
   caps.opMask = 0;
-  caps.defaultOp = GIN_OP_UNSET;
+  caps.defaultOp = ncclGinOpUnset;
   caps.isSignalOp = true;
   caps.allowAggregateRequests = false;
   caps.allowBidirectional = false;

src/device_api/gin/throughput/ginGetBW.cu

@@ -77,7 +77,7 @@ static void ginGetBWLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cud
 
 void ncclTestGinGetBWLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
                             const ginArgs_t* args, size_t numElems, int iters) {
-  if (args->ginRsm == GIN_RSM_CTA) {
+  if (args->ginRsm == ncclGinRsmCta) {
     ginGetBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, hBuf, stream, args, numElems, iters);
   } else {
     ginGetBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, hBuf, stream, args, numElems, iters);

src/device_api/gin/throughput/ginGetBW_main.cu

@@ -8,10 +8,10 @@ int main(int argc, char** argv) {
   ncclTestGinParseArgs(argc, argv, &args);
 
   ginTestCaps_t caps = {};
-  caps.defaultRsm = GIN_RSM_GPU;
+  caps.defaultRsm = ncclGinRsmGpu;
   caps.allowThreadRsm = false;
   caps.opMask = 0;
-  caps.defaultOp = GIN_OP_UNSET;
+  caps.defaultOp = ncclGinOpUnset;
   caps.isSignalOp = false;
   caps.allowAggregateRequests = true;
   caps.allowBidirectional = true;

src/device_api/gin/throughput/ginPutBW.cu

@@ -76,7 +76,7 @@ static void ginPutBWLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cud
 
 void ncclTestGinPutBWLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
                             const ginArgs_t* args, size_t numElems, int iters) {
-  if (args->ginRsm == GIN_RSM_CTA) {
+  if (args->ginRsm == ncclGinRsmCta) {
     ginPutBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, hBuf, stream, args, numElems, iters);
   } else {
     ginPutBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, hBuf, stream, args, numElems, iters);

src/device_api/gin/throughput/ginPutBW_main.cu

@@ -8,10 +8,10 @@ int main(int argc, char** argv) {
   ncclTestGinParseArgs(argc, argv, &args);
 
   ginTestCaps_t caps = {};
-  caps.defaultRsm = GIN_RSM_GPU;
+  caps.defaultRsm = ncclGinRsmGpu;
   caps.allowThreadRsm = false;
-  caps.opMask = GIN_OP_BIT(GIN_OP_PUT) | GIN_OP_BIT(GIN_OP_PUT_SIGNAL) | GIN_OP_BIT(GIN_OP_PUT_COUNTER);
-  caps.defaultOp = GIN_OP_PUT;
+  caps.opMask = GIN_OP_BIT(ncclGinOpPut) | GIN_OP_BIT(ncclGinOpPutSignal) | GIN_OP_BIT(ncclGinOpPutCount);
+  caps.defaultOp = ncclGinOpPut;
   caps.isSignalOp = false;
   caps.allowAggregateRequests = true;
   caps.allowBidirectional = true;
@@ -21,11 +21,11 @@ int main(int argc, char** argv) {
   ginBenchmark_t bench = {};
   bench.type = GIN_BENCHMARK_TYPE_THROUGHPUT;
   switch (args.ginOp) {
-  case GIN_OP_PUT_SIGNAL:
+  case ncclGinOpPutSignal:
     bench.run = ncclTestGinPutSignalBWLaunch;
     bench.name = "GIN PUT_SIGNAL BANDWIDTH/MESSAGE RATE";
     break;
-  case GIN_OP_PUT_COUNTER:
+  case ncclGinOpPutCount:
     bench.run = ncclTestGinPutCounterBWLaunch;
     bench.name = "GIN PUT_COUNTER BANDWIDTH/MESSAGE RATE";
     break;

src/device_api/gin/throughput/ginPutCounterBW.cu

@@ -81,7 +81,7 @@ static void ginPutCounterBWLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle hB
 
 void ncclTestGinPutCounterBWLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
                                    const ginArgs_t* args, size_t numElems, int iters) {
-  if (args->ginRsm == GIN_RSM_CTA) {
+  if (args->ginRsm == ncclGinRsmCta) {
     ginPutCounterBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, hBuf, stream, args, numElems, iters);
   } else {
     ginPutCounterBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, hBuf, stream, args, numElems, iters);

src/device_api/gin/throughput/ginPutSignalBW.cu

@@ -106,7 +106,7 @@ static void ginPutSignalBWLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle hBu
 
 void ncclTestGinPutSignalBWLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
                                   const ginArgs_t* args, size_t numElems, int iters) {
-  if (args->ginRsm == GIN_RSM_CTA) {
+  if (args->ginRsm == ncclGinRsmCta) {
     ginPutSignalBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, hBuf, stream, args, numElems, iters);
   } else {
     ginPutSignalBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, hBuf, stream, args, numElems, iters);

src/device_api/gin/throughput/ginSignalBW.cu

@@ -96,7 +96,7 @@ static void ginSignalBWLaunchRsm(ncclDevComm dcomm, ncclDevResourceHandle hBuf,
 
 void ncclTestGinSignalBWLaunch(ncclDevComm dcomm, ncclDevResourceHandle hBuf, cudaStream_t stream,
     const ginArgs_t* args, size_t numElems, int iters) {
-  if (args->ginRsm == GIN_RSM_CTA) {
+  if (args->ginRsm == ncclGinRsmCta) {
     ginSignalBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_CTA>(dcomm, hBuf, stream, args, iters);
   } else {
     ginSignalBWLaunchRsm<NCCL_GIN_RESOURCE_SHARING_GPU>(dcomm, hBuf, stream, args, iters);

src/device_api/gin/throughput/ginSignalBW_main.cu

@@ -11,10 +11,10 @@ int main(int argc, char** argv) {
   ncclTestGinParseArgs(argc, argv, &args);
 
   ginTestCaps_t caps = {};
-  caps.defaultRsm = GIN_RSM_GPU;
+  caps.defaultRsm = ncclGinRsmGpu;
   caps.allowThreadRsm = false;
   caps.opMask = 0;
-  caps.defaultOp = GIN_OP_UNSET;
+  caps.defaultOp = ncclGinOpUnset;
   caps.isSignalOp = true;
   caps.allowAggregateRequests = true;
   caps.allowBidirectional = true;

src/device_api/gin/utils/args.cc

@@ -69,38 +69,38 @@ static int parseIntArg(const char* str, const char* opt) {
   }
   return (int)val;
 }
-static ginRsm_t parseRsm(const char* str) {
-  if (strcmp(str, "thread") == 0) return GIN_RSM_THREAD;
-  if (strcmp(str, "cta") == 0) return GIN_RSM_CTA;
-  if (strcmp(str, "gpu") == 0) return GIN_RSM_GPU;
+static ncclGinRsm_t parseRsm(const char* str) {
+  if (strcmp(str, "thread") == 0) return ncclGinRsmThread;
+  if (strcmp(str, "cta") == 0) return ncclGinRsmCta;
+  if (strcmp(str, "gpu") == 0) return ncclGinRsmGpu;
   fprintf(stderr,
     "Error: invalid value for --gin_rsm: '%s' (expected thread, cta, or gpu)\n", str);
   exit(EXIT_FAILURE);
 }
 
-static ginOp_t parseOp(const char* str) {
-  if (strcmp(str, "put") == 0) return GIN_OP_PUT;
-  if (strcmp(str, "put_signal") == 0) return GIN_OP_PUT_SIGNAL;
-  if (strcmp(str, "put_counter") == 0) return GIN_OP_PUT_COUNTER;
+static ncclGinOp_t parseOp(const char* str) {
+  if (strcmp(str, "put") == 0) return ncclGinOpPut;
+  if (strcmp(str, "put_signal") == 0) return ncclGinOpPutSignal;
+  if (strcmp(str, "put_counter") == 0) return ncclGinOpPutCount;
   fprintf(stderr,
     "Error: invalid value for --gin_op: '%s' (expected put, put_signal, or put_counter)\n", str);
   exit(EXIT_FAILURE);
 }
 
-const char* ncclTestGinOpName(ginOp_t op) {
+const char* ncclTestGinOpName(ncclGinOp_t op) {
   switch (op) {
-  case GIN_OP_PUT: return "put";
-  case GIN_OP_PUT_SIGNAL: return "put_signal";
-  case GIN_OP_PUT_COUNTER: return "put_counter";
+  case ncclGinOpPut: return "put";
+  case ncclGinOpPutSignal: return "put_signal";
+  case ncclGinOpPutCount: return "put_counter";
   default: return "unset";
   }
 }
 
-const char* ncclTestGinRsmName(ginRsm_t rsm) {
+const char* ncclTestGinRsmName(ncclGinRsm_t rsm) {
   switch (rsm) {
-  case GIN_RSM_THREAD: return "thread";
-  case GIN_RSM_CTA: return "cta";
-  case GIN_RSM_GPU: return "gpu";
+  case ncclGinRsmThread: return "thread";
+  case ncclGinRsmCta: return "cta";
+  case ncclGinRsmGpu: return "gpu";
   default: return "unset";
   }
 }
@@ -119,8 +119,8 @@ void ncclTestGinParseArgs(int argc, char** argv, ginArgs_t* args) {
   args->ginAggregateRequests = false;
   args->ginBidirectional = false;
   args->queueDepth = 1024;
-  args->ginRsm = GIN_RSM_UNSET;
-  args->ginOp = GIN_OP_UNSET;
+  args->ginRsm = ncclGinRsmUnset;
+  args->ginOp = ncclGinOpUnset;
 
   static const struct option longOpts[] = {
     {"minbytes", required_argument, NULL, 'b'},
@@ -249,18 +249,18 @@ void ncclTestGinParseArgs(int argc, char** argv, ginArgs_t* args) {
 
 static void configureOp(ginArgs_t* args, const ginTestCaps_t* caps) {
   if (caps->opMask == 0) {
-    if (args->ginOp != GIN_OP_UNSET) {
+    if (args->ginOp != ncclGinOpUnset) {
       fprintf(stderr, "Error: --gin_op is only supported by the put benchmarks\n");
       exit(EXIT_FAILURE);
     }
     return;
   }
-  if (args->ginOp == GIN_OP_UNSET) args->ginOp = caps->defaultOp;
+  if (args->ginOp == ncclGinOpUnset) args->ginOp = caps->defaultOp;
   if ((caps->opMask & GIN_OP_BIT(args->ginOp)) == 0) {
     fprintf(stderr, "Error: --gin_op %s is not supported by this benchmark (accepted:",
             ncclTestGinOpName(args->ginOp));
-    for (int op = GIN_OP_PUT; op <= GIN_OP_PUT_COUNTER; op++) {
-      if (caps->opMask & GIN_OP_BIT(op)) fprintf(stderr, " %s", ncclTestGinOpName((ginOp_t)op));
+    for (int op = ncclGinOpPut; op <= ncclGinOpPutCount; op++) {
+      if (caps->opMask & GIN_OP_BIT(op)) fprintf(stderr, " %s", ncclTestGinOpName((ncclGinOp_t)op));
     }
     fprintf(stderr, ")\n");
     exit(EXIT_FAILURE);
@@ -268,15 +268,15 @@ static void configureOp(ginArgs_t* args, const ginTestCaps_t* caps) {
 }
 
 void ncclTestGinConfigureArgs(ginArgs_t* args, const ginTestCaps_t* caps) {
-  if (args->ginRsm == GIN_RSM_UNSET) args->ginRsm = caps->defaultRsm;
-  if (!caps->allowThreadRsm && args->ginRsm == GIN_RSM_THREAD) {
+  if (args->ginRsm == ncclGinRsmUnset) args->ginRsm = caps->defaultRsm;
+  if (!caps->allowThreadRsm && args->ginRsm == ncclGinRsmThread) {
     fprintf(stderr, "Error: --gin_rsm thread is not supported by throughput tests (expected cta or gpu)\n");
     exit(EXIT_FAILURE);
   }
 
   configureOp(args, caps);
 
-  const bool signalsInUse = caps->isSignalOp || args->ginOp == GIN_OP_PUT_SIGNAL;
+  const bool signalsInUse = caps->isSignalOp || args->ginOp == ncclGinOpPutSignal;
   if (!signalsInUse && args->ginStrongSignal) {
     fprintf(stderr,
       "Error: --gin_strong_signal is only supported by the signal tests and by put tests run with "

src/device_api/gin/utils/args.h

@@ -3,11 +3,11 @@
 #include <stddef.h>
 
 typedef enum {
-  GIN_RSM_UNSET = 0,
-  GIN_RSM_THREAD,
-  GIN_RSM_CTA,
-  GIN_RSM_GPU,
-} ginRsm_t;
+  ncclGinRsmUnset = 0,
+  ncclGinRsmThread,
+  ncclGinRsmCta,
+  ncclGinRsmGpu,
+} ncclGinRsm_t;
 
 /*
  * GIN operation under test, selected by --gin_op. A benchmark binary links one
@@ -15,11 +15,11 @@ typedef enum {
  * one, so e.g. the put binaries cover plain put, put+signal and put+counter.
  */
 typedef enum {
-  GIN_OP_UNSET = 0,
-  GIN_OP_PUT,
-  GIN_OP_PUT_SIGNAL,
-  GIN_OP_PUT_COUNTER,
-} ginOp_t;
+  ncclGinOpUnset = 0,
+  ncclGinOpPut,
+  ncclGinOpPutSignal,
+  ncclGinOpPutCount,
+} ncclGinOp_t;
 
 #define GIN_OP_BIT(op) (1u << (unsigned)(op))
 
@@ -37,8 +37,8 @@ typedef struct {
   bool ginAggregateRequests;
   bool ginBidirectional;
   int queueDepth;
-  ginRsm_t ginRsm;
-  ginOp_t ginOp;
+  ncclGinRsm_t ginRsm;
+  ncclGinOp_t ginOp;
 } ginArgs_t;
 
 /*
@@ -61,10 +61,10 @@ typedef struct {
  *                                 latency requires a single CTA and thread).
  */
 typedef struct {
-  ginRsm_t defaultRsm;
+  ncclGinRsm_t defaultRsm;
   bool allowThreadRsm;
   unsigned opMask;
-  ginOp_t defaultOp;
+  ncclGinOp_t defaultOp;
   bool isSignalOp;
   bool allowAggregateRequests;
   bool allowBidirectional;
@@ -73,5 +73,5 @@ typedef struct {
 
 void ncclTestGinParseArgs(int argc, char** argv, ginArgs_t* args);
 void ncclTestGinConfigureArgs(ginArgs_t* args, const ginTestCaps_t* caps);
-const char* ncclTestGinOpName(ginOp_t op);
-const char* ncclTestGinRsmName(ginRsm_t rsm);
+const char* ncclTestGinOpName(ncclGinOp_t op);
+const char* ncclTestGinRsmName(ncclGinRsm_t rsm);