diff --git a/src/include/devcomm.h b/src/include/devcomm.h
index 963dab5428..194df1b7af 100644
--- a/src/include/devcomm.h
+++ b/src/include/devcomm.h
@@ -220,6 +220,16 @@ struct ncclChannel {
       int collFifoTail; // Only used by CPU
 
       uint32_t* sync;
+#ifdef ENABLE_PROFILING
+      struct timeval tv_send[NCCL_STEPS];
+      struct timeval tv_recv[NCCL_STEPS];
+      uint64_t send_byte;
+      uint64_t recv_byte;
+      uint64_t send_us;
+      uint64_t send_wait_us;
+      uint64_t recv_us;
+      uint64_t recv_flush_us;
+#endif
     };
     int data[0x80];
   };
diff --git a/src/init.cc b/src/init.cc
index ac5d80730f..bea2defcd2 100644
--- a/src/init.cc
+++ b/src/init.cc
@@ -241,6 +241,17 @@ static ncclResult_t commFree(ncclComm_t comm) {
     (prof->recvCopySend_cycle) ? (double)prof->recvCopySend_byte*comm->nChannels/((double)prof->recvCopySend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0);
   free(prof);
   CUDACHECK(hipFree(comm->hostDevComm.devProf));
+
+  for (int channel=0; channel<comm->p2pnChannels; channel++) {
+    if (comm->channels[channel].send_byte) INFO(NCCL_INIT, "# [%03d:%02d] Proxy Send %6.2f GB/s (%ld bytes %ld+%ld us)",
+      comm->rank, channel, (comm->channels[channel].send_us+comm->channels[channel].send_wait_us) ?
+      (double)comm->channels[channel].send_byte/(comm->channels[channel].send_us+comm->channels[channel].send_wait_us)/1000.0 : 0,
+      comm->channels[channel].send_byte, comm->channels[channel].send_wait_us, comm->channels[channel].send_us);
+    if (comm->channels[channel].recv_byte) INFO(NCCL_INIT, "# [%03d:%02d] Proxy Recv %6.2f GB/s (%ld bytes %ld+%ld us)",
+      comm->rank, channel, (comm->channels[channel].recv_us+comm->channels[channel].recv_flush_us) ?
+      (double)comm->channels[channel].recv_byte/(comm->channels[channel].recv_us+comm->channels[channel].recv_flush_us)/1000.0 : 0,
+      comm->channels[channel].recv_byte, comm->channels[channel].recv_us, comm->channels[channel].recv_flush_us);
+  }
 #endif
 
 #ifdef ENABLE_COLLTRACE
diff --git a/src/transport/net.cc b/src/transport/net.cc
index 8478b9345d..850912332c 100644
--- a/src/transport/net.cc
+++ b/src/transport/net.cc
@@ -8,6 +8,7 @@
 #include "comm.h"
 #include "net.h"
 #include "graph.h"
+#include <sys/time.h>
 
 struct netConnectInfo {
   ncclNetHandle_t netHandle;
@@ -302,8 +303,20 @@ ncclResult_t netSendProxy(struct ncclProxyArgs* args) {
         } else if (args->tail < LOAD(recvTail)) {
           // Send through network
           if (LOAD(sizesFifo+buffSlot) != -1) {
+#ifdef ENABLE_PROFILING
+            if (args->channel->tv_send[buffSlot].tv_sec == 0 && args->channel->tv_send[buffSlot].tv_usec == 0)
+              gettimeofday(&args->channel->tv_send[buffSlot], NULL);
+#endif
             NCCLCHECK(ncclNetIsend(resources->netSendComm, localBuff+buffSlot*stepSize, sizesFifo[buffSlot], mhandle, args->requests+buffSlot));
             if (args->requests[buffSlot] != NULL) {
+#ifdef ENABLE_PROFILING
+              struct timeval tv;
+              gettimeofday(&tv, NULL);
+              if (args->opCount > 0) args->channel->send_wait_us += ((uint64_t)(tv.tv_sec - args->channel->tv_send[buffSlot].tv_sec)*1000*1000 + tv.tv_usec - args->channel->tv_send[buffSlot].tv_usec);
+              if (args->opCount > 0) args->channel->send_byte += LOAD(sizesFifo+buffSlot);
+              args->channel->tv_send[buffSlot].tv_sec = tv.tv_sec;
+              args->channel->tv_send[buffSlot].tv_usec = tv.tv_usec;
+#endif
               STORE(sizesFifo+buffSlot, -1);
               // Make sure size is reset to zero before we update the head.
               __sync_synchronize();
@@ -318,6 +331,13 @@ ncclResult_t netSendProxy(struct ncclProxyArgs* args) {
         int buffSlot = args->head%NCCL_STEPS;
         NCCLCHECK(ncclNetTest(args->requests[buffSlot], &done, NULL));
         if (done) {
+#ifdef ENABLE_PROFILING
+          struct timeval tv;
+          gettimeofday(&tv, NULL);
+          if (args->opCount > 0) args->channel->send_us += ((uint64_t)(tv.tv_sec - args->channel->tv_send[buffSlot].tv_sec)*1000*1000 + tv.tv_usec - args->channel->tv_send[buffSlot].tv_usec);
+          args->channel->tv_send[buffSlot].tv_sec = 0;
+          args->channel->tv_send[buffSlot].tv_usec = 0;
+#endif
           args->head += args->sliceSteps;
           STORE(&resources->sendMem->head, args->head);
           args->idle = 0;
@@ -357,6 +377,9 @@ ncclResult_t netRecvProxy(struct ncclProxyArgs* args) {
       if ((args->tail < args->head + NCCL_STEPS) && (args->tail < LOAD(sendHead) + NCCL_STEPS) && (args->tail < args->end)) {
         int buffSlot = args->tail%NCCL_STEPS;
         int sliceSize = stepSize * args->sliceSteps;
+#ifdef ENABLE_PROFILING
+        gettimeofday(&args->channel->tv_recv[buffSlot], NULL);
+#endif
         NCCLCHECK(ncclNetIrecv(resources->netRecvComm, localBuff+buffSlot*stepSize, sliceSize, mhandle, args->requests+buffSlot));
         if (args->requests[buffSlot] != NULL) {
           args->tail += args->sliceSteps;
@@ -370,7 +393,19 @@ ncclResult_t netRecvProxy(struct ncclProxyArgs* args) {
         if (done) {
           args->head += args->sliceSteps;
           if (args->protocol == NCCL_PROTO_SIMPLE) {
+#ifdef ENABLE_PROFILING
+            struct timeval tv;
+            gettimeofday(&tv, NULL);
+            if (args->opCount > 0) args->channel->recv_byte += size;
+            if (args->opCount > 0) args->channel->recv_us += ((uint64_t)(tv.tv_sec - args->channel->tv_recv[buffSlot].tv_sec)*1000*1000 + tv.tv_usec - args->channel->tv_recv[buffSlot].tv_usec);
+            args->channel->tv_recv[buffSlot].tv_sec = tv.tv_sec;
+            args->channel->tv_recv[buffSlot].tv_usec = tv.tv_usec;
+#endif
             if (resources->useGdr) NCCLCHECK(ncclNetFlush(resources->netRecvComm, localBuff+buffSlot*stepSize, size, mhandle));
+#ifdef ENABLE_PROFILING
+            gettimeofday(&tv, NULL);
+            if (args->opCount > 0) args->channel->recv_flush_us += ((uint64_t)(tv.tv_sec - args->channel->tv_recv[buffSlot].tv_sec)*1000*1000 + tv.tv_usec - args->channel->tv_recv[buffSlot].tv_usec);
+#endif
             STORE(&resources->recvMem->tail, args->head);
           }
           args->idle = 0;