From 7a6c6927aee5ed362fbc17ebb521f04dc030f9ef Mon Sep 17 00:00:00 2001 From: Wenkai Du <43822138+wenkaidu@users.noreply.github.com> Date: Fri, 3 Jun 2022 07:05:13 -0700 Subject: [PATCH] Enable timing profile option (#558) --- src/collectives/device/all_reduce.h | 2 ++ src/collectives/device/primitives.h | 4 ++++ src/include/devcomm.h | 1 + src/init.cc | 21 ++++++++++++++++++++- 4 files changed, 27 insertions(+), 1 deletion(-) diff --git a/src/collectives/device/all_reduce.h b/src/collectives/device/all_reduce.h index 2ee9dc0f22..0529653248 100644 --- a/src/collectives/device/all_reduce.h +++ b/src/collectives/device/all_reduce.h @@ -37,8 +37,10 @@ namespace { minChunkSize = nthreads*(Proto::calcBytePerGrain()/sizeof(T))/2; } + INIT_COUNTER; Primitives, 0, Proto, 0> prims (tid, nthreads, &ring->prev, &ring->next, args->sendbuff, args->recvbuff, args->redOpArg, args->connIndex << 16); + ACCUMULATE_PRIM_COUNTER(prim); for (ssize_t gridOffset = 0; gridOffset < size; gridOffset += loopSize) { ssize_t realChunkSize; diff --git a/src/collectives/device/primitives.h b/src/collectives/device/primitives.h index ee8f06a569..ce97c48232 100644 --- a/src/collectives/device/primitives.h +++ b/src/collectives/device/primitives.h @@ -163,9 +163,13 @@ struct PrimitivesWithoutDirect { if (tid == 0) { struct ncclProfElem *elem = devProf.elems+args->opCount%PROFILE_NUM_ITEMS; elem->elem[blockIdx.x].prim##_cycle += (__builtin_amdgcn_s_memrealtime() - t0 \ + ws - elem->elem[blockIdx.x].wait_cycle); \ elem->elem[blockIdx.x].prim##_byte += nelem * sizeof(T); elem->elem[blockIdx.x].opCount = args->opCount;} +#define ACCUMULATE_PRIM_COUNTER(prim) \ + if (tid == 0) { struct ncclProfElem *elem = devProf.elems+args->opCount%PROFILE_NUM_ITEMS; elem->elem[blockIdx.x].prim##_cycle += (__builtin_amdgcn_s_memrealtime() - t0 \ + + ws - elem->elem[blockIdx.x].wait_cycle); elem->elem[blockIdx.x].opCount = args->opCount;} #else #define INIT_COUNTER #define ACCUMULATE_COUNTER(prim) +#define ACCUMULATE_PRIM_COUNTER(prim) #endif #endif diff --git a/src/include/devcomm.h b/src/include/devcomm.h index 22cbd94c5d..0c7fb3dfbe 100644 --- a/src/include/devcomm.h +++ b/src/include/devcomm.h @@ -327,6 +327,7 @@ struct ncclProfElem { uint64_t total_cycle; uint64_t wait_cycle; // total wait cycle // primtive cycles + uint64_t prim_cycle; uint64_t send_cycle; uint64_t directSend_cycle; uint64_t recv_cycle; diff --git a/src/init.cc b/src/init.cc index 275620044a..112757fe09 100644 --- a/src/init.cc +++ b/src/init.cc @@ -229,6 +229,7 @@ void *ncclCommThreadMain(void *arg) { #endif #undef NCCL_NO_OPTIMIZE +#define PROFILE_USE_TIME static ncclResult_t commFree(ncclComm_t comm) { if (comm == NULL) @@ -255,8 +256,12 @@ static ncclResult_t commFree(ncclComm_t comm) { CUDACHECK(hipMemcpy(prof.elems, comm->hostDevComm.devProf.elems, sizeof(struct ncclProfElem)*PROFILE_NUM_ITEMS, hipMemcpyDeviceToHost)); #define VEGA_GPU_RTC_FREQUENCY 2.5E7 if (comm->rank == 0) { - INFO(NCCL_INIT, "# %7s %4s %6s %6s %6s %6s %7s %6s %6s %6s %6s %6s", "Rank:Ch", "opCt", "total", " wait", "send", "rcRdS", "dRcRdCS", "dRcCS", "dRc", "cS", "rc", "rcCS"); + INFO(NCCL_INIT, "# %7s %4s %6s %6s %6s %6s %6s %7s %6s %6s %6s %6s %6s", "Rank:Ch", "opCt", "total", " prim", " wait", "send", "rcRdS", "dRcRdCS", "dRcCS", "dRc", "cS", "rc", "rcCS"); +#ifdef PROFILE_USE_TIME + INFO(NCCL_INIT, "# %7s %4s %6s %6s %6s %6s %6s %7s %6s %6s %6s %6s %6s", "", "", " (us)", " (us)", " (us)", " (us)", " (us)", " (us)", " (us)", " (us)", " (us)", " (us)", " (us)"); +#else INFO(NCCL_INIT, "# %7s %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)"); +#endif } for (int i = 1; i < PROFILE_NUM_ITEMS; i++) { int valid = 0; @@ -265,6 +270,19 @@ static ncclResult_t commFree(ncclComm_t comm) { if (elem->elem[chan].opCount == 0) continue; valid++; +#ifdef PROFILE_USE_TIME + INFO(NCCL_INIT, "# [%02d:%02d] %04d %6.2f %6.2f %6.2f %6.2f %6.2f %7.2f %6.2f %6.2f %6.2f %6.2f %6.2f", + comm->rank, chan, (uint32_t)elem->elem[chan].opCount, (double)elem->elem[chan].total_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E6, + (double)elem->elem[chan].prim_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E6, (double)elem->elem[chan].wait_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E6, + (elem->elem[chan].send_cycle) ? ((double)elem->elem[chan].send_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E6) : 0, + (elem->elem[chan].recvReduceSend_cycle) ? ((double)elem->elem[chan].recvReduceSend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E6) : 0, + (elem->elem[chan].directRecvReduceCopySend_cycle) ? ((double)elem->elem[chan].directRecvReduceCopySend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E6) : 0, + (elem->elem[chan].directRecvCopySend_cycle) ? ((double)elem->elem[chan].directRecvCopySend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E6) : 0, + (elem->elem[chan].directRecv_cycle) ? ((double)elem->elem[chan].directRecv_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E6) : 0, + (elem->elem[chan].copySend_cycle) ? ((double)elem->elem[chan].copySend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E6) : 0, + (elem->elem[chan].recv_cycle) ? ((double)elem->elem[chan].recv_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E6) : 0, + (elem->elem[chan].recvCopySend_cycle) ? ((double)elem->elem[chan].recvCopySend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E6) : 0); +#else INFO(NCCL_INIT, "# [%02d:%02d] %04d %6.4f %6.4f %6.2f %6.2f %7.2f %6.2f %6.2f %6.2f %6.2f %6.2f", comm->rank, chan, (uint32_t)elem->elem[chan].opCount, (double)elem->elem[chan].total_cycle/VEGA_GPU_RTC_FREQUENCY, (double)elem->elem[chan].wait_cycle/VEGA_GPU_RTC_FREQUENCY, @@ -276,6 +294,7 @@ static ncclResult_t commFree(ncclComm_t comm) { (elem->elem[chan].copySend_cycle) ? (double)elem->elem[chan].copySend_byte/((double)elem->elem[chan].copySend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0, (elem->elem[chan].recv_cycle) ? (double)elem->elem[chan].recv_byte/((double)elem->elem[chan].recv_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0, (elem->elem[chan].recvCopySend_cycle) ? (double)elem->elem[chan].recvCopySend_byte/((double)elem->elem[chan].recvCopySend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0); +#endif } if (valid == 0) break;