AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Enable timing profiling mode (#447)

Browse Source 
[ROCm/rccl commit: ec36c4c326]
This commit is contained in:
Wenkai Du
2021-10-27 08:21:48 -07:00
committed by GitHub
parent a6feafd5dc
commit 747216e2b2
7 changed files with 71 additions and 12 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl/CMakeLists.txt
+5
View File
@@ -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)
projects/rccl/src/collectives/device/all_reduce.h
+1 -1
View File
@@ -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
}
projects/rccl/src/collectives/device/broadcast.h
+1 -1
View File
@@ -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
}
}
projects/rccl/src/collectives/device/primitives.h
+9 -2
View File
@@ -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)
projects/rccl/src/collectives/device/prims_simple.h
+12 -8
View File
@@ -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<DirectRecv, DirectSend, Recv, Send, Src, Dst>(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<RedOp>::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;
projects/rccl/src/init.cc
+35
View File
@@ -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; chan<comm->nChannels; 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; channel<std::max(comm->nChannels, 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);
projects/rccl/src/transport/net.cc
+8
View File
@@ -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 ++;
}
}