diff --git a/src/collectives/device/prims_simple.h b/src/collectives/device/prims_simple.h
index 32eb2a13a0..acf0c95b36 100644
--- a/src/collectives/device/prims_simple.h
+++ b/src/collectives/device/prims_simple.h
@@ -86,6 +86,9 @@ class Primitives<
   inline __device__ void waitPeer(intptr_t dstIx, intptr_t remoteOutIx, int offset, int nelts) {
     if (flags & (Recv*RoleWaitRecv | Send*RoleWaitSend)) {
       bool const isSendNotRecv = (Send && Recv) ? (flags & RoleWaitSend) : Send;
+#ifdef ENABLE_PROFILING
+      uint64_t t0 = __builtin_amdgcn_s_memrealtime();
+#endif
       int spins = 0;
       while (connStepCache + (isSendNotRecv ? NCCL_STEPS : 0) < step + StepPerSlice) {
         __builtin_amdgcn_s_sleep(8);
@@ -107,6 +110,12 @@ class Primitives<
       else
         ptrs[index] = connEltsFifo + (step%NCCL_STEPS)*stepSize;
       step += StepPerSlice;
+#ifdef ENABLE_PROFILING
+      if (isSendNotRecv)
+        ncclShmem->comm.devProf->elems[blockIdx.x].wait_send_cycle += (__builtin_amdgcn_s_memrealtime() - t0);
+      else
+        ncclShmem->comm.devProf->elems[blockIdx.x].wait_recv_cycle += (__builtin_amdgcn_s_memrealtime() - t0);
+#endif
     }
   }
 
diff --git a/src/graph/topo.cc b/src/graph/topo.cc
index d859dd7c4e..0e72141769 100644
--- a/src/graph/topo.cc
+++ b/src/graph/topo.cc
@@ -379,7 +379,7 @@ ncclResult_t ncclTopoAddGpu(struct ncclXmlNode* xmlGpu, struct ncclTopoSystem* s
 }
 
 struct kvDict kvDictPciClass[] = { { "0x060400", PCI }, { "0x068000", NVS }, { "0x068001", CPU }, { "0x030200", GPU }, { "0x030000", GPU }, { "0x038000", GPU }, { "0x020700", NIC }, { "0x020000", NIC }, { NULL, PCI /* Default fallback value */ } };
-struct kvDict kvDictPciGen[] = { { "2.5 GT/s", 15 }, { "5 GT/s", 30 }, { "8 GT/s", 60 }, { "16 GT/s", 120 }, { NULL, 60 /* Default fallback */ } }; // x100 Mbps per lane
+struct kvDict kvDictPciGen[] = { { "2.5 GT/s", 15 }, { "5 GT/s", 30 }, { "8 GT/s", 60 }, { "16 GT/s", 120 }, { "32 GT/s", 240 }, { "8.0 GT/s", 60 }, { "16.0 GT/s", 120 }, { "32.0 GT/s", 240 }, { NULL, 60 /* Default fallback */ } }; // x100 Mbps per lane
 ncclResult_t ncclTopoAddPci(struct ncclXmlNode* xmlPci, struct ncclTopoSystem* system, struct ncclTopoNode* parent) {
   const char* str;
 
diff --git a/src/include/devcomm.h b/src/include/devcomm.h
index ae95f19cd0..08398a7e3a 100644
--- a/src/include/devcomm.h
+++ b/src/include/devcomm.h
@@ -288,6 +288,8 @@ struct ncclProfElem {
     struct {
       uint64_t total_cycle;
       uint64_t wait_cycle;      // total wait cycle
+      uint64_t wait_send_cycle;
+      uint64_t wait_recv_cycle;
       // primtive cycles
       uint64_t send_cycle;
       uint64_t directSend_cycle;
diff --git a/src/init.cc b/src/init.cc
index a4f3ce3a10..d2dd9ce27d 100644
--- a/src/init.cc
+++ b/src/init.cc
@@ -312,7 +312,7 @@ 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 total_cycle = 0, wait_cycle = 0, send_cycle = 0, directSend_cycle = 0, recv_cycle = 0, \
+  uint64_t total_cycle = 0, wait_cycle = 0, wait_send_cycle = 0, wait_recv_cycle = 0, send_cycle = 0, directSend_cycle = 0, recv_cycle = 0, \
     directRecv_cycle = 0, copySend_cycle = 0, directCopySend_cycle = 0, recvCopySend_cycle = 0, directRecvCopySend_cycle = 0, \
     recvReduceCopy_cycle = 0, recvReduceSend_cycle = 0, recvReduceCopySend_cycle = 0, directRecvReduceCopySend_cycle = 0, \
     send_byte = 0, directSend_byte = 0, recv_byte = 0, directRecv_byte = 0, copySend_byte = 0, directCopySend_byte = 0, \
@@ -321,6 +321,8 @@ static ncclResult_t commFree(ncclComm_t comm) {
   for (int chan=0; chan<comm->nChannels; chan++) {
     total_cycle += prof->elems[chan].total_cycle;
     wait_cycle += prof->elems[chan].wait_cycle;
+    wait_send_cycle += prof->elems[chan].wait_send_cycle;
+    wait_recv_cycle += prof->elems[chan].wait_recv_cycle;
     send_cycle += prof->elems[chan].send_cycle;
     directSend_cycle += prof->elems[chan].directSend_cycle;
     recv_cycle += prof->elems[chan].recv_cycle;
@@ -348,12 +350,14 @@ static ncclResult_t commFree(ncclComm_t comm) {
   }
   #define VEGA_GPU_RTC_FREQUENCY 2.5E7
   if (comm->rank == 0) {
-    INFO(NCCL_INIT, "# %4s %6s %6s %6s %6s %7s %6s %6s %6s %6s %6s", "Rank", "total", "  wait", "send", "rcRdS", "dRcRdCS", "dRcCS", "dRc", "cS", "rc", "rcCS");
-    INFO(NCCL_INIT, "# %4s %6s %6s %6s %6s %7s %6s %6s %6s %6s %6s", "", "(s)", "(s)", "(GB/s)", "(GB/s)", "(GB/s)", "(GB/s)", "(GB/s)", "(GB/s)", "(GB/s)", "(GB/s)");
+    INFO(NCCL_INIT, "# %4s %6s %6s %6s %6s %6s %6s %7s %6s %6s %6s %6s %6s", "Rank", "total", "  wait", "w_send", "w_recv", "send", "rcRdS", "dRcRdCS", "dRcCS", "dRc", "cS", "rc", "rcCS");
+    INFO(NCCL_INIT, "# %4s %6s %6s %6s %6s %6s %6s %7s %6s %6s %6s %6s %6s", "", "(s)", "(s)", "(GB/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.2f %6.2f %7.2f %6.2f %6.2f %6.2f %6.2f %6.2f",
+  INFO(NCCL_INIT, "# %4d %6.4f %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)total_cycle/VEGA_GPU_RTC_FREQUENCY/comm->nChannels,
     (double)wait_cycle/VEGA_GPU_RTC_FREQUENCY/comm->nChannels,
+    (double)wait_send_cycle/VEGA_GPU_RTC_FREQUENCY/comm->nChannels,
+    (double)wait_recv_cycle/VEGA_GPU_RTC_FREQUENCY/comm->nChannels,
     (send_cycle) ? (double)send_byte*comm->nChannels/((double)send_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0,
     (recvReduceSend_cycle) ? (double)recvReduceSend_byte*comm->nChannels/((double)recvReduceSend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0,
     (directRecvReduceCopySend_cycle) ? (double)directRecvReduceCopySend_byte*comm->nChannels/((double)directRecvReduceCopySend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0,