From cac33a8c2f757e20b3e9219abb3210cff3245364 Mon Sep 17 00:00:00 2001
From: Wenkai Du <wenkai.du@amd.com>
Date: Thu, 3 Apr 2025 17:31:54 -0500
Subject: [PATCH] Automatically set in-place option from out-of-place (#123)

---
 README.md     |  3 ++-
 src/common.cu | 26 ++++++++++++++++++--------
 src/common.h  |  1 +
 3 files changed, 21 insertions(+), 9 deletions(-)

diff --git a/README.md b/README.md
index 4fdb74ae0d..f0f6347c39 100644
--- a/README.md
+++ b/README.md
@@ -139,8 +139,9 @@ All tests support the same set of arguments :
   * `-G,--hipgraph <num graph launches>` Capture iterations as a HIP graph and then replay specified number of times. Default : 0.
   * `-C,--report_cputime <0/1>]` Report CPU time instead of latency. Default : 0.
   * `-R,--local_register <1/0>` enable local buffer registration on send/recv buffers. Default : 0.
-  * `-T,--timeout <time in seconds>` timeout each test after specified number of seconds. Default : disabled.
+  * `-T,--timeout <time in seconds>` timeout each test after specified number of seconds. Default: disabled.
   * `-F,--cache_flush <cache flush after every -F iteration>` Enable cache flush after every -F iteration. Default : 0 (No cache flush).
+  * `-O,--out_of_place <0=in-place only, 1=out-of-place only>`. Default: both.
   * `-q,--delay <delay>` Delay between out-of-place and in-place runs (in microseconds). Default: 10.
 * Parsing RCCL-Tests output
   * `-Z,--output_format <csv|json>` Parse RCCL-Tests output as a CSV or JSON. Default : disabled.
diff --git a/src/common.cu b/src/common.cu
index 1ae2a788d4..cb546db011 100644
--- a/src/common.cu
+++ b/src/common.cu
@@ -111,6 +111,7 @@ static int average = 1;
 static int numDevices = 1;
 static int delay_inout_place = 0;
 static int enable_out_of_place = 1;
+static int enable_in_place = 1;
 static int enable_cache_flush = 0;
 static int enable_rotating_tensor = 0;
 #if NCCL_VERSION_CODE >= NCCL_VERSION(2,19,0)
@@ -410,7 +411,7 @@ testResult_t CheckData(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t
 
   int64_t *wrongPerGpu = nullptr;
   CUDACHECK(hipHostMalloc((void**)&wrongPerGpu, args->nGpus*sizeof(int64_t), cudaHostAllocMapped));
-  
+
   for (int i=0; i<args->nGpus; i++) {
     int rank = ((args->proc*args->nThreads + args->thread)*args->nGpus + i);
     CUDACHECK(cudaSetDevice(args->gpus[i]));
@@ -450,14 +451,14 @@ testResult_t CheckData(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t
   if (args->reportErrors && *wrongElts) args->errors[0]++;
   return testSuccess;
 }
-    
+
 testResult_t testStreamSynchronize(int ngpus, cudaStream_t* streams, ncclComm_t* comms) {
   cudaError_t cudaErr;
   int remaining = ngpus;
   int* done = (int*)malloc(sizeof(int)*ngpus);
   memset(done, 0, sizeof(int)*ngpus);
   timer tim;
-  
+
   while (remaining) {
    int idle = 1;
    for (int i=0; i<ngpus; i++) {
@@ -522,7 +523,7 @@ testResult_t startColl(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t
     size_t steps = totalnbytes ? args->maxbytes / totalnbytes : 1;
     shift = totalnbytes * (iter % steps);
   }
-  
+
   if (args->nGpus > 1) NCCLCHECK(ncclGroupStart());
   for (int i = 0; i < args->nGpus; i++) {
 #ifndef NCCL_MAJOR
@@ -912,7 +913,8 @@ testResult_t TimeTest(struct threadArgs* args, ncclDataType_t type, const char*
         TESTCHECK(BenchTime(args, type, op, root, 0));
         usleep(delay_inout_place);
       }
-      TESTCHECK(BenchTime(args, type, op, root, 1));
+        if (enable_in_place)
+        TESTCHECK(BenchTime(args, type, op, root, 1));
       PRINT("\n");
     }
     --repeat;
@@ -1206,10 +1208,11 @@ int main(int argc, char* argv[]) {
 	break;
       case 'O':
         enable_out_of_place = strtol(optarg, NULL, 0);
+        enable_in_place = enable_out_of_place ? 0 : 1;
         break;
       case 'q':
         delay_inout_place = (int)strtol(optarg, NULL, 10);
-	break;
+      	break;
       case 'F':
         enable_cache_flush = strtol(optarg, NULL, 0);
         if (enable_cache_flush > 0) {
@@ -1500,14 +1503,20 @@ testResult_t run() {
 
   const char* timeStr = report_cputime ? "cputime" : "time";
   PRINT("#\n");
-  if (enable_out_of_place) {
+  if (enable_out_of_place && enable_in_place) {
   	PRINT("# %10s  %12s  %8s  %6s  %6s           out-of-place                       in-place          \n", "", "", "", "", "");
   	PRINT("# %10s  %12s  %8s  %6s  %6s  %7s  %6s  %6s %6s  %7s  %6s  %6s %6s\n", "size", "count", "type", "redop", "root",
       	timeStr, "algbw", "busbw", "#wrong", timeStr, "algbw", "busbw", "#wrong");
   	PRINT("# %10s  %12s  %8s  %6s  %6s  %7s  %6s  %6s  %5s  %7s  %6s  %6s  %5s\n", "(B)", "(elements)", "", "", "",
       	"(us)", "(GB/s)", "(GB/s)", "", "(us)", "(GB/s)", "(GB/s)", "");
+  } else if (enable_out_of_place) {
+	  PRINT("# %10s  %12s  %8s  %6s  %6s           out-of-place      \n", "", "", "", "", "");
+        PRINT("# %10s  %12s  %8s  %6s  %6s  %7s  %6s  %6s %6s\n", "size", "count", "type", "redop", "root",
+        timeStr, "algbw", "busbw", "#wrong");
+        PRINT("# %10s  %12s  %8s  %6s  %6s  %7s  %6s  %6s  %5s\n", "(B)", "(elements)", "", "", "",
+        "(us)", "(GB/s)", "(GB/s)", "");
   } else {
-	PRINT("# %10s  %12s  %8s  %6s  %6s           in-place          \n", "", "", "", "", "");
+    PRINT("# %10s  %12s  %8s  %6s  %6s           in-place          \n", "", "", "", "", "");
         PRINT("# %10s  %12s  %8s  %6s  %6s  %7s  %6s  %6s %6s\n", "size", "count", "type", "redop", "root",
         timeStr, "algbw", "busbw", "#wrong");
         PRINT("# %10s  %12s  %8s  %6s  %6s  %7s  %6s  %6s  %5s\n", "(B)", "(elements)", "", "", "",
@@ -1539,6 +1548,7 @@ testResult_t run() {
     threads[t].args.comms=comms+t*nGpus;
     threads[t].args.streams=streams.data()+t*nGpus;
     threads[t].args.enable_out_of_place=enable_out_of_place;
+    threads[t].args.enable_in_place=enable_in_place;
     threads[t].args.enable_cache_flush = enable_cache_flush;
     threads[t].args.enable_rotating_tensor = enable_rotating_tensor;
     threads[t].args.errors=errors.data()+t;
diff --git a/src/common.h b/src/common.h
index 961059c578..1b368cd28a 100644
--- a/src/common.h
+++ b/src/common.h
@@ -157,6 +157,7 @@ struct threadArgs {
   int* gpus;
   int localRank;
   int enable_out_of_place;
+  int enable_in_place;
   int enable_cache_flush;
   int enable_rotating_tensor;
   void** sendbuffs;