From 5d93f83ae34b1c590eca680f01f97be2f8220b71 Mon Sep 17 00:00:00 2001
From: Wenkai Du <wenkai.du@amd.com>
Date: Mon, 15 Jun 2020 20:51:36 +0000
Subject: [PATCH] Limit network profiling support to simple protocol and avoid
 overflow

[ROCm/rccl commit: 95b8f70d15f87a430696e52f966c069eff38dd22]
---
 projects/rccl/src/include/devcomm.h |  2 +-
 projects/rccl/src/transport/net.cc  | 34 +++++++++++++++++++----------
 2 files changed, 23 insertions(+), 13 deletions(-)

diff --git a/projects/rccl/src/include/devcomm.h b/projects/rccl/src/include/devcomm.h
index dd362af16d..e5c5f43727 100644
--- a/projects/rccl/src/include/devcomm.h
+++ b/projects/rccl/src/include/devcomm.h
@@ -222,7 +222,7 @@ struct ncclChannel {
       uint32_t* sync;
 #ifdef ENABLE_PROFILING
       struct timeval tvs;
-      int sizes;
+      uint64_t sizes;
       int active_req;
       uint64_t send_byte;
       uint64_t recv_byte;
diff --git a/projects/rccl/src/transport/net.cc b/projects/rccl/src/transport/net.cc
index 87803d574c..7a9bf960a5 100644
--- a/projects/rccl/src/transport/net.cc
+++ b/projects/rccl/src/transport/net.cc
@@ -340,12 +340,17 @@ ncclResult_t netSendProxy(struct ncclProxyArgs* args) {
         NCCLCHECK(ncclNetTest(args->requests[buffSlot], &done, NULL));
         if (done) {
 #ifdef ENABLE_PROFILING
-          args->channel->active_req --;
-          if (args->channel->active_req == 0) {
-            struct timeval tv;
-            gettimeofday(&tv, NULL);
-            args->channel->bw_cumulative += (float)args->channel->sizes/((tv.tv_sec - args->channel->tvs.tv_sec)*1000*1000 + tv.tv_usec - args->channel->tvs.tv_usec)/1000.0;
-            args->channel->bw_count ++;
+          if (args->protocol == NCCL_PROTO_SIMPLE) {
+            args->channel->active_req --;
+            if (args->channel->active_req == 0) {
+              struct timeval tv;
+              gettimeofday(&tv, NULL);
+              float delta = (tv.tv_sec - args->channel->tvs.tv_sec)*1E6 + tv.tv_usec - args->channel->tvs.tv_usec;
+              if (delta) {
+                args->channel->bw_cumulative += (float)args->channel->sizes/delta/1E3;
+                args->channel->bw_count ++;
+              }
+            }
           }
 #endif
           args->head += args->sliceSteps;
@@ -390,11 +395,13 @@ ncclResult_t netRecvProxy(struct ncclProxyArgs* args) {
         NCCLCHECK(ncclNetIrecv(resources->netRecvComm, localBuff+buffSlot*stepSize, sliceSize, mhandle, args->requests+buffSlot));
         if (args->requests[buffSlot] != NULL) {
 #ifdef ENABLE_PROFILING
-          if (args->channel->active_req == 0) {
-            gettimeofday(&args->channel->tvs, NULL);
-            args->channel->sizes = 0;
+          if (args->protocol == NCCL_PROTO_SIMPLE) {
+            if (args->channel->active_req == 0) {
+              gettimeofday(&args->channel->tvs, NULL);
+              args->channel->sizes = 0;
+            }
+            args->channel->active_req ++;
           }
-          args->channel->active_req ++;
 #endif
           args->tail += args->sliceSteps;
           args->idle = 0;
@@ -414,8 +421,11 @@ ncclResult_t netRecvProxy(struct ncclProxyArgs* args) {
           if (args->channel->active_req == 0) {
             struct timeval tv;
             gettimeofday(&tv, NULL);
-            args->channel->bw_cumulative += (float)args->channel->sizes/((tv.tv_sec - args->channel->tvs.tv_sec)*1000*1000 + tv.tv_usec - args->channel->tvs.tv_usec)/1000.0;
-            args->channel->bw_count ++;
+            float delta = (tv.tv_sec - args->channel->tvs.tv_sec)*1E6 + tv.tv_usec - args->channel->tvs.tv_usec;
+            if (delta) {
+              args->channel->bw_cumulative += (float)args->channel->sizes/delta/1E3;
+              args->channel->bw_count ++;
+            }
           }
 #endif
             if (resources->useGdr) NCCLCHECK(ncclNetFlush(resources->netRecvComm, localBuff+buffSlot*stepSize, size, mhandle));