From 747216e2b2879f074061fec4bfec1df95397ce9f Mon Sep 17 00:00:00 2001 From: Wenkai Du <43822138+wenkaidu@users.noreply.github.com> Date: Wed, 27 Oct 2021 08:21:48 -0700 Subject: [PATCH] Enable timing profiling mode (#447) [ROCm/rccl commit: ec36c4c32628cf8f17917788a75ade0d8303078d] --- projects/rccl/CMakeLists.txt | 5 +++ .../rccl/src/collectives/device/all_reduce.h | 2 +- .../rccl/src/collectives/device/broadcast.h | 2 +- .../rccl/src/collectives/device/primitives.h | 11 ++++-- .../src/collectives/device/prims_simple.h | 20 ++++++----- projects/rccl/src/init.cc | 35 +++++++++++++++++++ projects/rccl/src/transport/net.cc | 8 +++++ 7 files changed, 71 insertions(+), 12 deletions(-) diff --git a/projects/rccl/CMakeLists.txt b/projects/rccl/CMakeLists.txt index 0bf53d3823..9718d218b7 100644 --- a/projects/rccl/CMakeLists.txt +++ b/projects/rccl/CMakeLists.txt @@ -199,6 +199,11 @@ if(PROFILE) add_definitions(-DENABLE_PROFILING) endif() +if(TIMING_PROFILE) + add_definitions(-DENABLE_PROFILING) + add_definitions(-DENABLE_TIMING_PROFILE) +endif() + set(COLLTRACE 1 CACHE BOOL "Collective Trace Option") if(COLLTRACE) add_definitions(-DENABLE_COLLTRACE) diff --git a/projects/rccl/src/collectives/device/all_reduce.h b/projects/rccl/src/collectives/device/all_reduce.h index 673cf6390e..97f347a887 100644 --- a/projects/rccl/src/collectives/device/all_reduce.h +++ b/projects/rccl/src/collectives/device/all_reduce.h @@ -110,7 +110,7 @@ namespace { ACCUMULATE_COUNTER(directRecv); } #ifdef ENABLE_PROFILING - if (tid == 0) devProf->elems[blockIdx.x].total_cycle += (__builtin_amdgcn_s_memrealtime() - clk); + if (tid == 0 && args->op.opCount) devProf->elems[blockIdx.x].total_cycle += (__builtin_amdgcn_s_memrealtime() - clk); #endif } diff --git a/projects/rccl/src/collectives/device/broadcast.h b/projects/rccl/src/collectives/device/broadcast.h index 20cc37539e..c318afd93c 100644 --- a/projects/rccl/src/collectives/device/broadcast.h +++ b/projects/rccl/src/collectives/device/broadcast.h @@ -70,7 +70,7 @@ namespace { } } #ifdef ENABLE_PROFILING - if (tid == 0) devProf->elems[blockIdx.x].total_cycle += (__builtin_amdgcn_s_memrealtime() - clk); + if (tid == 0 && args->op.opCount) devProf->elems[blockIdx.x].total_cycle += (__builtin_amdgcn_s_memrealtime() - clk); #endif } } diff --git a/projects/rccl/src/collectives/device/primitives.h b/projects/rccl/src/collectives/device/primitives.h index 2b500e230c..9fa1a427c0 100644 --- a/projects/rccl/src/collectives/device/primitives.h +++ b/projects/rccl/src/collectives/device/primitives.h @@ -157,13 +157,20 @@ struct PrimitivesWithoutDirect { #include "prims_ll128.h" #ifdef ENABLE_PROFILING +#ifdef ENABLE_TIMING_PROFILE +#define INIT_COUNTER \ + if (tid == 0) { t0 = __builtin_amdgcn_s_memrealtime(); } +#define ACCUMULATE_COUNTER(prim) \ + if (tid == 0 && args->op.opCount) { devProf->elems[blockIdx.x].prim##_cycle += (__builtin_amdgcn_s_memrealtime() - t0); \ + devProf->elems[blockIdx.x].prim##_byte += nelem * sizeof(T); } +#else #define INIT_COUNTER \ if (tid == 0) { t0 = __builtin_amdgcn_s_memrealtime(); ws = devProf->elems[blockIdx.x].wait_cycle; } - #define ACCUMULATE_COUNTER(prim) \ - if (tid == 0) { devProf->elems[blockIdx.x].prim##_cycle += (__builtin_amdgcn_s_memrealtime() - t0 \ + if (tid == 0 && args->op.opCount) { devProf->elems[blockIdx.x].prim##_cycle += (__builtin_amdgcn_s_memrealtime() - t0 \ + ws - devProf->elems[blockIdx.x].wait_cycle); \ devProf->elems[blockIdx.x].prim##_byte += nelem * sizeof(T); } +#endif #else #define INIT_COUNTER #define ACCUMULATE_COUNTER(prim) diff --git a/projects/rccl/src/collectives/device/prims_simple.h b/projects/rccl/src/collectives/device/prims_simple.h index acf0c95b36..92f2d29c93 100644 --- a/projects/rccl/src/collectives/device/prims_simple.h +++ b/projects/rccl/src/collectives/device/prims_simple.h @@ -46,6 +46,7 @@ class Primitives< uint64_t* barriers; uint64_t* barrier_next; const int connIndex; + const uint64_t opCount; // Don't use barrier 0 as it's used by the final sync inline __device__ void barrier() { @@ -86,7 +87,7 @@ 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 +#if defined(ENABLE_PROFILING) && !defined(ENABLE_TIMING_PROFILE) uint64_t t0 = __builtin_amdgcn_s_memrealtime(); #endif int spins = 0; @@ -110,11 +111,13 @@ 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); +#if defined(ENABLE_PROFILING) && !defined(ENABLE_TIMING_PROFILE) + if (opCount) { + 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 } } @@ -180,7 +183,7 @@ class Primitives< waitPeer(dstIx, remoteOutIx, offset, sliceSize); subBarrier(); #ifdef ENABLE_PROFILING - if (tid == 0) ncclShmem->comm.devProf->elems[blockIdx.x].wait_cycle += (__builtin_amdgcn_s_memrealtime() - t0); + if (tid == 0 && opCount) ncclShmem->comm.devProf->elems[blockIdx.x].wait_cycle += (__builtin_amdgcn_s_memrealtime() - t0); #endif if (DirectRecv && ncclShmem->groups[group].srcs[0] == ncclShmem->groups[group].dsts[0]) { // We can only have one direct receive. Since srcs[0] == dstPtr+offset, skip one copy @@ -341,7 +344,8 @@ class Primitives< stepSize(ncclShmem->comm.buffSizes[NCCL_PROTO_SIMPLE]/NCCL_STEPS/sizeof(T)), redOp(FuncTraits::make(ncclShmem->comm.nRanks)), connIndex((NCCL_MAX_DIRECT_ARITY==Fan::MaxSend || NCCL_MAX_DIRECT_ARITY==Fan::MaxRecv)?(group/2):connIndex), - barriers(&ncclShmem->groups[group].barrier), barrier_next(ncclShmem->groups[group].barrier_next) { + barriers(&ncclShmem->groups[group].barrier), barrier_next(ncclShmem->groups[group].barrier_next), + opCount(ncclShmem->work.elems[0].op.opCount) { // For send operations, we need an extra warp to overlap the threadfence and the copy this->nthreads = nthreads; diff --git a/projects/rccl/src/init.cc b/projects/rccl/src/init.cc index d2dd9ce27d..78ff291fd3 100644 --- a/projects/rccl/src/init.cc +++ b/projects/rccl/src/init.cc @@ -310,6 +310,40 @@ static ncclResult_t commFree(ncclComm_t comm) { free(comm->asyncOps); #ifdef ENABLE_PROFILING +#ifdef ENABLE_TIMING_PROFILE + struct ncclProf* prof = (struct ncclProf*)malloc(sizeof(struct ncclProf)); + CUDACHECK(hipMemcpy(prof, comm->hostDevComm.devProf, sizeof(struct ncclProf), hipMemcpyDeviceToHost)); + #define VEGA_GPU_RTC_FREQUENCY 2.5E7 + if (comm->rank == 0) { + INFO(NCCL_INIT, "# %8s %7s %7s %7s %7s %7s %7s %7s %7s %7s", "Rank:Ch", "total", "send", "rcRdS", "dRcRdCS", "dRcCS", "dRc", "cS", "rc", "rcCS"); + INFO(NCCL_INIT, "# %8s %7s %7s %7s %7s %7s %7s %7s %7s %7s", "", "(ms)", "(ms)", "(ms)", "(ms)", "(ms)", "(ms)", "(ms)", "(ms)", "(ms)"); + } + for (int chan=0; channChannels; chan++) { + INFO(NCCL_INIT, "# [%03d:%02d] %7.3f %7.3f %7.3f %7.3f %7.3f %7.3f %7.3f %7.3f %7.3f", + comm->rank, chan, (double)prof->elems[chan].total_cycle/VEGA_GPU_RTC_FREQUENCY*1000.0, + (double)prof->elems[chan].send_cycle/VEGA_GPU_RTC_FREQUENCY*1000.0, + (double)prof->elems[chan].recvReduceSend_cycle/VEGA_GPU_RTC_FREQUENCY*1000.0, + (double)prof->elems[chan].directRecvReduceCopySend_cycle/VEGA_GPU_RTC_FREQUENCY*1000.0, + (double)prof->elems[chan].directRecvCopySend_cycle/VEGA_GPU_RTC_FREQUENCY*1000.0, + (double)prof->elems[chan].directRecv_cycle/VEGA_GPU_RTC_FREQUENCY*1000.0, + (double)prof->elems[chan].copySend_cycle/VEGA_GPU_RTC_FREQUENCY*1000.0, + (double)prof->elems[chan].recv_cycle/VEGA_GPU_RTC_FREQUENCY*1000.0, + (double)prof->elems[chan].recvCopySend_cycle/VEGA_GPU_RTC_FREQUENCY*1000.0); + } + free(prof); + CUDACHECK(hipFree(comm->hostDevComm.devProf)); + + for (int channel=0; channelnChannels, comm->p2pnChannels); channel++) { + if (comm->channels[channel].send_byte) INFO(NCCL_INIT, "# [%03d:%02d] Proxy Send %7.3f ms (%ld bytes %d measurements)", + comm->rank, channel, (comm->channels[channel].bw_count) ? + (float)comm->channels[channel].bw_cumulative/comm->channels[channel].bw_count : 0, + comm->channels[channel].send_byte, comm->channels[channel].bw_count); + if (comm->channels[channel].recv_byte) INFO(NCCL_INIT, "# [%03d:%02d] Proxy Recv %7.3f ms (%ld bytes %d measurements)", + comm->rank, channel, (comm->channels[channel].bw_count) ? + (float)comm->channels[channel].bw_cumulative/comm->channels[channel].bw_count : 0, + comm->channels[channel].recv_byte, comm->channels[channel].bw_count); + } +#else 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, wait_send_cycle = 0, wait_recv_cycle = 0, send_cycle = 0, directSend_cycle = 0, recv_cycle = 0, \ @@ -380,6 +414,7 @@ static ncclResult_t commFree(ncclComm_t comm) { comm->channels[channel].recv_byte, comm->channels[channel].bw_count); } #endif +#endif #ifdef ENABLE_COLLTRACE STORE(&comm->hostDevComm.collTraceExit, 1); diff --git a/projects/rccl/src/transport/net.cc b/projects/rccl/src/transport/net.cc index 701e091945..be24232ca9 100644 --- a/projects/rccl/src/transport/net.cc +++ b/projects/rccl/src/transport/net.cc @@ -431,7 +431,11 @@ ncclResult_t netSendProxy(struct ncclProxyArgs* args) { gettimeofday(&tv, NULL); float delta = (tv.tv_sec - sub->channel->tvs.tv_sec)*1E6 + tv.tv_usec - sub->channel->tvs.tv_usec; if (delta) { +#ifdef ENABLE_TIMING_PROFILE + sub->channel->bw_cumulative += (float)delta/1E3; +#else sub->channel->bw_cumulative += (float)sub->channel->sizes/delta/1E3; +#endif sub->channel->bw_count ++; } } @@ -526,7 +530,11 @@ ncclResult_t netRecvProxy(struct ncclProxyArgs* args) { gettimeofday(&tv, NULL); float delta = (tv.tv_sec - sub->channel->tvs.tv_sec)*1E6 + tv.tv_usec - sub->channel->tvs.tv_usec; if (delta) { +#ifdef ENABLE_TIMING_PROFILE + sub->channel->bw_cumulative += (float)delta/1E3; +#else sub->channel->bw_cumulative += (float)sub->channel->sizes/delta/1E3; +#endif sub->channel->bw_count ++; } }