AI/rocm-systems
2
0
Vork 0
Code Kwesties Pull-aanvragen Actions Paketten Projecten Publicaties Wiki Activiteit

Calculate and use total wait cycles for RCCL profiling

Bladeren bron 
[ROCm/rccl commit: c9aa11928a]
This commit is contained in:
Wenkai Du
2020-05-22 15:03:52 -07:00
bovenliggende 15cd2aff3c
commit bc6840f937
5 gewijzigde bestanden met toevoegingen van 20 en 25 verwijderingen
Download Patch-bestand Download Diff-bestand Expand all files Collapse all files
projects/rccl/src/collectives/device/all_reduce.h
+2 -2
Bestand weergeven
@@ -26,7 +26,7 @@ __device__ void ncclAllReduceRingKernel(struct CollectiveArgs* args) {
const ssize_t size = args->coll.count;
#ifdef ENABLE_PROFILING
auto devProf = comm->devProf;
uint64_t clk, t0 = 0ULL, ws, wr;
uint64_t clk, t0 = 0ULL, ws;
if (tid == 0) clk = __rtc64();
#endif
@@ -99,7 +99,7 @@ __device__ void ncclAllReduceRingKernel(struct CollectiveArgs* args) {
ACCUMULATE_COUNTER(directRecv);
}
#ifdef ENABLE_PROFILING
if (tid == 0) __atomic_fetch_add(&(devProf->total_cycle), __rtc64() - clk, __ATOMIC_SEQ_CST);
if (tid == 0 && args->opCount > 0) __atomic_fetch_add(&(devProf->total_cycle), __rtc64() - clk, __ATOMIC_SEQ_CST);
#endif
}
projects/rccl/src/collectives/device/broadcast.h
+2 -2
Bestand weergeven
@@ -28,7 +28,7 @@ __device__ void ncclBroadcastRingKernel(struct CollectiveArgs* args) {
const int root = args->coll.root;
#ifdef ENABLE_PROFILING
auto devProf = comm->devProf;
uint64_t clk, t0 = 0ULL, ws, wr;
uint64_t clk, t0 = 0ULL, ws;
if (tid == 0) clk = __rtc64();
#endif
@@ -66,7 +66,7 @@ __device__ void ncclBroadcastRingKernel(struct CollectiveArgs* args) {
}
}
#ifdef ENABLE_PROFILING
if (tid == 0) __atomic_fetch_add(&(devProf->total_cycle), __rtc64() - clk, __ATOMIC_SEQ_CST);
if (tid == 0 && args->opCount > 0) __atomic_fetch_add(&(devProf->total_cycle), __rtc64() - clk, __ATOMIC_SEQ_CST);
#endif
}
projects/rccl/src/collectives/device/primitives.h
+10 -15
Bestand weergeven
@@ -135,17 +135,10 @@ class ncclPrimitives {
spins = 0;
mismatch = 0;
if (s.sendConnHeadPtr) {
#ifdef ENABLE_PROFILING
auto devProf = comm->devProf;
uint64_t t0 = __rtc64();
#endif
while (s.sendConnHeadCache + NCCL_STEPS < s.sendConnHead + SLICESTEPS) {
s.sendConnHeadCache = LOAD(s.sendConnHeadPtr);
if (checkAbort(wid, 1)) break;
}
#ifdef ENABLE_PROFILING
__atomic_fetch_add(&devProf->wait_send_cycle[blockIdx.x], __rtc64() - t0, __ATOMIC_SEQ_CST);
#endif
if (s.sendConnFifoPtr) {
STORE(s.sendConnFifoPtr+s.sendConnHead%NCCL_STEPS, nbytes);
}
@@ -158,7 +151,6 @@ class ncclPrimitives {
mismatch = 0;
if (r.recvConnTailPtr) {
#ifdef ENABLE_PROFILING
auto devProf = comm->devProf;
uint64_t t0 = __rtc64();
#endif
while (r.recvConnTailCache < r.recvConnTail + SLICESTEPS) {
@@ -166,7 +158,7 @@ class ncclPrimitives {
if (checkAbort(wid, 0)) break;
}
#ifdef ENABLE_PROFILING
__atomic_fetch_add(&devProf->wait_recv_cycle[blockIdx.x], __rtc64() - t0, __ATOMIC_SEQ_CST);
if (opCount > 0) __atomic_fetch_add(&comm->devProf->wait_recv_cycle[blockIdx.x], __rtc64() - t0, __ATOMIC_SEQ_CST);
#endif
r.recvConnTail += SLICESTEPS;
}
@@ -249,10 +241,16 @@ inline __device__ int directSendInc(int i, int directInc, int sliceInc) {
#pragma unroll
for (int slice=0; slice<SLICESPERCHUNK; ++slice) {
int realSize = max(0, min(dataSize, nelem-offset));
#ifdef ENABLE_PROFILING
uint64_t t0 = __rtc64();
#endif
if (SEND) waitSend(realSize*sizeof(T));
if (RECV) waitRecv();
if (realSize > 0) {
barrier();
#ifdef ENABLE_PROFILING
if (tid == 0 && opCount > 0) __atomic_fetch_add(&comm->devProf->wait_cycle[blockIdx.x], __rtc64() - t0, __ATOMIC_SEQ_CST);
#endif
#if defined(RCCL_USE_DIRECT_BUFFER)
if (DIRECTRECV && r.recvDirectBuff[0]) {
// We can only have one direct receive. Since srcs[0] == dstPtr+offset, skip one copy
@@ -457,14 +455,11 @@ inline __device__ int directSendInc(int i, int directInc, int sliceInc) {
#ifdef ENABLE_PROFILING
#define INIT_COUNTER \
if (tid == 0) { t0 = __rtc64(); ws = LOAD(&(devProf->wait_send_cycle[blockIdx.x])); \
wr = LOAD(&(devProf->wait_recv_cycle[blockIdx.x])); }
if (tid == 0) { t0 = __rtc64(); ws = LOAD(&(devProf->wait_cycle[blockIdx.x])); }
#define ACCUMULATE_COUNTER(prim) \
if (tid == 0) { __atomic_fetch_add(&(devProf->prim##_cycle), __rtc64() - t0 \
+ ws - LOAD(&(devProf->wait_send_cycle[blockIdx.x])) \
+ wr - LOAD(&(devProf->wait_recv_cycle[blockIdx.x])), \
__ATOMIC_SEQ_CST); \
if (tid == 0 && args->opCount > 0) { __atomic_fetch_add(&(devProf->prim##_cycle), __rtc64() - t0 \
+ ws - LOAD(&(devProf->wait_cycle[blockIdx.x])), __ATOMIC_SEQ_CST); \
__atomic_fetch_add(&(devProf->prim##_byte), nelem * sizeof(T), __ATOMIC_SEQ_CST); }
#else
#define INIT_COUNTER
projects/rccl/src/include/devcomm.h
+2 -2
Bestand weergeven
@@ -232,8 +232,8 @@ struct ncclProf {
union {
struct {
uint64_t total_cycle;
uint64_t wait_send_cycle[MAXCHANNELS];
uint64_t wait_recv_cycle[MAXCHANNELS];
uint64_t wait_cycle[MAXCHANNELS]; // total wait cycle
uint64_t wait_recv_cycle[MAXCHANNELS]; // recv wait cycle
// primtive cycles
uint64_t send_cycle;
uint64_t directSend_cycle;
projects/rccl/src/init.cc
+4 -4
Bestand weergeven
@@ -217,19 +217,19 @@ 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 wait_send_cycle = 0, wait_recv_cycle = 0;
uint64_t wait_cycle = 0, wait_recv_cycle = 0;
for (int chan=0; chan<comm->nChannels; chan++) {
wait_send_cycle += prof->wait_send_cycle[chan];
wait_cycle += prof->wait_cycle[chan];
wait_recv_cycle += prof->wait_recv_cycle[chan];
}
#define VEGA_GPU_RTC_FREQUENCY 2.5E7
if (comm->rank == 0) {
INFO(NCCL_INIT, "# %4s %6s %6s %6s %6s %6s %7s %6s %6s %6s %6s %6s", "Rank", "total", "w_send", "w_recv", "send", "rcRdS", "dRcRdCS", "dRcCS", "dRc", "cS", "rc", "rcCS");
INFO(NCCL_INIT, "# %4s %6s %6s %6s %6s %6s %7s %6s %6s %6s %6s %6s", "Rank", "total", " wait", "w_recv", "send", "rcRdS", "dRcRdCS", "dRcCS", "dRc", "cS", "rc", "rcCS");
INFO(NCCL_INIT, "# %4s %6s %6s %6s %6s %6s %7s %6s %6s %6s %6s %6s", "", "(s)", "(s)", "(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.4f %6.2f %6.2f %7.2f %6.2f %6.2f %6.2f %6.2f %6.2f",
comm->rank, (double)prof->total_cycle/VEGA_GPU_RTC_FREQUENCY/comm->nChannels,
(double)wait_send_cycle/VEGA_GPU_RTC_FREQUENCY/comm->nChannels,
(double)wait_cycle/VEGA_GPU_RTC_FREQUENCY/comm->nChannels,
(double)wait_recv_cycle/VEGA_GPU_RTC_FREQUENCY/comm->nChannels,
(prof->send_cycle) ? (double)prof->send_byte*comm->nChannels/((double)prof->send_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0,
(prof->recvReduceSend_cycle) ? (double)prof->recvReduceSend_byte*comm->nChannels/((double)prof->recvReduceSend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0,