diff --git a/projects/rccl/src/collectives/device/all_reduce.h b/projects/rccl/src/collectives/device/all_reduce.h
index 14aaba580c..0ea08cdaa7 100644
--- a/projects/rccl/src/collectives/device/all_reduce.h
+++ b/projects/rccl/src/collectives/device/all_reduce.h
@@ -26,7 +26,7 @@ __device__ void ncclAllReduceRingKernel(struct CollectiveArgs* args) {
   const ssize_t size = args->coll.count;
 #ifdef ENABLE_PROFILING
   auto devProf = comm->devProf;
-  uint64_t clk, t0 = 0ULL, ws, wr;
+  uint64_t clk, t0 = 0ULL, ws;
   if (tid == 0) clk = __rtc64();
 #endif
 
@@ -99,7 +99,7 @@ __device__ void ncclAllReduceRingKernel(struct CollectiveArgs* args) {
     ACCUMULATE_COUNTER(directRecv);
   }
 #ifdef ENABLE_PROFILING
-  if (tid == 0) __atomic_fetch_add(&(devProf->total_cycle), __rtc64() - clk, __ATOMIC_SEQ_CST);
+  if (tid == 0 && args->opCount > 0) __atomic_fetch_add(&(devProf->total_cycle), __rtc64() - clk, __ATOMIC_SEQ_CST);
 #endif
 }
 
diff --git a/projects/rccl/src/collectives/device/broadcast.h b/projects/rccl/src/collectives/device/broadcast.h
index 235f9102a8..6c186b20f0 100644
--- a/projects/rccl/src/collectives/device/broadcast.h
+++ b/projects/rccl/src/collectives/device/broadcast.h
@@ -28,7 +28,7 @@ __device__ void ncclBroadcastRingKernel(struct CollectiveArgs* args) {
   const int root = args->coll.root;
 #ifdef ENABLE_PROFILING
   auto devProf = comm->devProf;
-  uint64_t clk, t0 = 0ULL, ws, wr;
+  uint64_t clk, t0 = 0ULL, ws;
   if (tid == 0) clk = __rtc64();
 #endif
 
@@ -66,7 +66,7 @@ __device__ void ncclBroadcastRingKernel(struct CollectiveArgs* args) {
     }
   }
 #ifdef ENABLE_PROFILING
-  if (tid == 0) __atomic_fetch_add(&(devProf->total_cycle), __rtc64() - clk, __ATOMIC_SEQ_CST);
+  if (tid == 0 && args->opCount > 0) __atomic_fetch_add(&(devProf->total_cycle), __rtc64() - clk, __ATOMIC_SEQ_CST);
 #endif
 }
 
diff --git a/projects/rccl/src/collectives/device/primitives.h b/projects/rccl/src/collectives/device/primitives.h
index f8d22db5b8..023004cb3e 100644
--- a/projects/rccl/src/collectives/device/primitives.h
+++ b/projects/rccl/src/collectives/device/primitives.h
@@ -135,17 +135,10 @@ class ncclPrimitives {
     spins = 0;
     mismatch = 0;
     if (s.sendConnHeadPtr) {
-#ifdef ENABLE_PROFILING
-      auto devProf = comm->devProf;
-      uint64_t t0 = __rtc64();
-#endif
       while (s.sendConnHeadCache + NCCL_STEPS < s.sendConnHead + SLICESTEPS) {
         s.sendConnHeadCache = LOAD(s.sendConnHeadPtr);
         if (checkAbort(wid, 1)) break;
       }
-#ifdef ENABLE_PROFILING
-      __atomic_fetch_add(&devProf->wait_send_cycle[blockIdx.x], __rtc64() - t0, __ATOMIC_SEQ_CST);
-#endif
       if (s.sendConnFifoPtr) {
         STORE(s.sendConnFifoPtr+s.sendConnHead%NCCL_STEPS, nbytes);
       }
@@ -158,7 +151,6 @@ class ncclPrimitives {
     mismatch = 0;
     if (r.recvConnTailPtr) {
 #ifdef ENABLE_PROFILING
-      auto devProf = comm->devProf;
       uint64_t t0 = __rtc64();
 #endif
       while (r.recvConnTailCache < r.recvConnTail + SLICESTEPS) {
@@ -166,7 +158,7 @@ class ncclPrimitives {
         if (checkAbort(wid, 0)) break;
       }
 #ifdef ENABLE_PROFILING
-      __atomic_fetch_add(&devProf->wait_recv_cycle[blockIdx.x], __rtc64() - t0, __ATOMIC_SEQ_CST);
+      if (opCount > 0) __atomic_fetch_add(&comm->devProf->wait_recv_cycle[blockIdx.x], __rtc64() - t0, __ATOMIC_SEQ_CST);
 #endif
       r.recvConnTail += SLICESTEPS;
     }
@@ -249,10 +241,16 @@ inline __device__ int directSendInc(int i, int directInc, int sliceInc) {
     #pragma unroll
     for (int slice=0; slice<SLICESPERCHUNK; ++slice) {
       int realSize = max(0, min(dataSize, nelem-offset));
+#ifdef ENABLE_PROFILING
+      uint64_t t0 = __rtc64();
+#endif
       if (SEND) waitSend(realSize*sizeof(T));
       if (RECV) waitRecv();
       if (realSize > 0) {
         barrier();
+#ifdef ENABLE_PROFILING
+        if (tid == 0  && opCount > 0) __atomic_fetch_add(&comm->devProf->wait_cycle[blockIdx.x], __rtc64() - t0, __ATOMIC_SEQ_CST);
+#endif
 #if defined(RCCL_USE_DIRECT_BUFFER)
         if (DIRECTRECV && r.recvDirectBuff[0]) {
           // We can only have one direct receive. Since srcs[0] == dstPtr+offset, skip one copy
@@ -457,14 +455,11 @@ inline __device__ int directSendInc(int i, int directInc, int sliceInc) {
 
 #ifdef ENABLE_PROFILING
 #define INIT_COUNTER \
-  if (tid == 0) { t0 = __rtc64(); ws = LOAD(&(devProf->wait_send_cycle[blockIdx.x])); \
-    wr = LOAD(&(devProf->wait_recv_cycle[blockIdx.x])); }
+  if (tid == 0) { t0 = __rtc64(); ws = LOAD(&(devProf->wait_cycle[blockIdx.x])); }
 
 #define ACCUMULATE_COUNTER(prim) \
-  if (tid == 0) { __atomic_fetch_add(&(devProf->prim##_cycle), __rtc64() - t0 \
-    + ws - LOAD(&(devProf->wait_send_cycle[blockIdx.x])) \
-    + wr - LOAD(&(devProf->wait_recv_cycle[blockIdx.x])), \
-    __ATOMIC_SEQ_CST); \
+  if (tid == 0 && args->opCount > 0) { __atomic_fetch_add(&(devProf->prim##_cycle), __rtc64() - t0 \
+    + ws - LOAD(&(devProf->wait_cycle[blockIdx.x])), __ATOMIC_SEQ_CST); \
     __atomic_fetch_add(&(devProf->prim##_byte), nelem * sizeof(T), __ATOMIC_SEQ_CST); }
 #else
 #define INIT_COUNTER
diff --git a/projects/rccl/src/include/devcomm.h b/projects/rccl/src/include/devcomm.h
index eabf190b96..963dab5428 100644
--- a/projects/rccl/src/include/devcomm.h
+++ b/projects/rccl/src/include/devcomm.h
@@ -232,8 +232,8 @@ struct ncclProf {
   union {
     struct {
       uint64_t total_cycle;
-      uint64_t wait_send_cycle[MAXCHANNELS];
-      uint64_t wait_recv_cycle[MAXCHANNELS];
+      uint64_t wait_cycle[MAXCHANNELS];      // total wait cycle
+      uint64_t wait_recv_cycle[MAXCHANNELS]; // recv wait cycle
       // primtive cycles
       uint64_t send_cycle;
       uint64_t directSend_cycle;
diff --git a/projects/rccl/src/init.cc b/projects/rccl/src/init.cc
index ac760bba73..ac5d80730f 100644
--- a/projects/rccl/src/init.cc
+++ b/projects/rccl/src/init.cc
@@ -217,19 +217,19 @@ static ncclResult_t commFree(ncclComm_t comm) {
 #ifdef ENABLE_PROFILING
   struct ncclProf* prof = (struct ncclProf*)malloc(sizeof(struct ncclProf));
   CUDACHECK(hipMemcpy(prof, comm->hostDevComm.devProf, sizeof(struct ncclProf), hipMemcpyDeviceToHost));
-  uint64_t wait_send_cycle = 0, wait_recv_cycle = 0;
+  uint64_t wait_cycle = 0, wait_recv_cycle = 0;
   for (int chan=0; chan<comm->nChannels; chan++) {
-    wait_send_cycle += prof->wait_send_cycle[chan];
+    wait_cycle += prof->wait_cycle[chan];
     wait_recv_cycle += prof->wait_recv_cycle[chan];
   }
   #define VEGA_GPU_RTC_FREQUENCY 2.5E7
   if (comm->rank == 0) {
-    INFO(NCCL_INIT, "# %4s %6s %6s %6s %6s %6s %7s %6s %6s %6s %6s %6s", "Rank", "total", "w_send", "w_recv", "send", "rcRdS", "dRcRdCS", "dRcCS", "dRc", "cS", "rc", "rcCS");
+    INFO(NCCL_INIT, "# %4s %6s %6s %6s %6s %6s %7s %6s %6s %6s %6s %6s", "Rank", "total", "  wait", "w_recv", "send", "rcRdS", "dRcRdCS", "dRcCS", "dRc", "cS", "rc", "rcCS");
     INFO(NCCL_INIT, "# %4s %6s %6s %6s %6s %6s %7s %6s %6s %6s %6s %6s", "", "(s)", "(s)", "(s)", "(GB/s)", "(GB/s)", "(GB/s)", "(GB/s)", "(GB/s)", "(GB/s)", "(GB/s)", "(GB/s)", "(GB/s)");
   }
   INFO(NCCL_INIT, "# %4d %6.4f %6.4f %6.4f %6.2f %6.2f %7.2f %6.2f %6.2f %6.2f %6.2f %6.2f",
     comm->rank, (double)prof->total_cycle/VEGA_GPU_RTC_FREQUENCY/comm->nChannels,
-    (double)wait_send_cycle/VEGA_GPU_RTC_FREQUENCY/comm->nChannels,
+    (double)wait_cycle/VEGA_GPU_RTC_FREQUENCY/comm->nChannels,
     (double)wait_recv_cycle/VEGA_GPU_RTC_FREQUENCY/comm->nChannels,
     (prof->send_cycle) ? (double)prof->send_byte*comm->nChannels/((double)prof->send_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0,
     (prof->recvReduceSend_cycle) ? (double)prof->recvReduceSend_byte*comm->nChannels/((double)prof->recvReduceSend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0,