AI/rocm-systems
2
0
Förgrening 0
Kod Ärenden Pull-förfrågningar Actions Packages Projekt Släpp Wiki Aktiviteter

Add network proxy profiling support

Bläddra i källkod
This commit is contained in:
Wenkai Du
2020-05-28 00:15:47 +00:00
förälder c9aa11928a
incheckning 812543104d
3 ändrade filer med 56 tillägg och 0 borttagningar
Download Patch File Download Diff File Expand all files Collapse all files
src/include/devcomm.h
+10
Visa fil
@@ -220,6 +220,16 @@ struct ncclChannel {
int collFifoTail; // Only used by CPU
uint32_t* sync;
#ifdef ENABLE_PROFILING
struct timeval tv_send[NCCL_STEPS];
struct timeval tv_recv[NCCL_STEPS];
uint64_t send_byte;
uint64_t recv_byte;
uint64_t send_us;
uint64_t send_wait_us;
uint64_t recv_us;
uint64_t recv_flush_us;
#endif
};
int data[0x80];
};
src/init.cc
+11
Visa fil
@@ -241,6 +241,17 @@ static ncclResult_t commFree(ncclComm_t comm) {
(prof->recvCopySend_cycle) ? (double)prof->recvCopySend_byte*comm->nChannels/((double)prof->recvCopySend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0);
free(prof);
CUDACHECK(hipFree(comm->hostDevComm.devProf));
for (int channel=0; channel<comm->p2pnChannels; channel++) {
if (comm->channels[channel].send_byte) INFO(NCCL_INIT, "# [%03d:%02d] Proxy Send %6.2f GB/s (%ld bytes %ld+%ld us)",
comm->rank, channel, (comm->channels[channel].send_us+comm->channels[channel].send_wait_us) ?
(double)comm->channels[channel].send_byte/(comm->channels[channel].send_us+comm->channels[channel].send_wait_us)/1000.0 : 0,
comm->channels[channel].send_byte, comm->channels[channel].send_wait_us, comm->channels[channel].send_us);
if (comm->channels[channel].recv_byte) INFO(NCCL_INIT, "# [%03d:%02d] Proxy Recv %6.2f GB/s (%ld bytes %ld+%ld us)",
comm->rank, channel, (comm->channels[channel].recv_us+comm->channels[channel].recv_flush_us) ?
(double)comm->channels[channel].recv_byte/(comm->channels[channel].recv_us+comm->channels[channel].recv_flush_us)/1000.0 : 0,
comm->channels[channel].recv_byte, comm->channels[channel].recv_us, comm->channels[channel].recv_flush_us);
}
#endif
#ifdef ENABLE_COLLTRACE
src/transport/net.cc
+35
Visa fil
@@ -8,6 +8,7 @@
#include "comm.h"
#include "net.h"
#include "graph.h"
#include <sys/time.h>
struct netConnectInfo {
ncclNetHandle_t netHandle;
@@ -302,8 +303,20 @@ ncclResult_t netSendProxy(struct ncclProxyArgs* args) {
} else if (args->tail < LOAD(recvTail)) {
// Send through network
if (LOAD(sizesFifo+buffSlot) != -1) {
#ifdef ENABLE_PROFILING
if (args->channel->tv_send[buffSlot].tv_sec == 0 && args->channel->tv_send[buffSlot].tv_usec == 0)
gettimeofday(&args->channel->tv_send[buffSlot], NULL);
#endif
NCCLCHECK(ncclNetIsend(resources->netSendComm, localBuff+buffSlot*stepSize, sizesFifo[buffSlot], mhandle, args->requests+buffSlot));
if (args->requests[buffSlot] != NULL) {
#ifdef ENABLE_PROFILING
struct timeval tv;
gettimeofday(&tv, NULL);
if (args->opCount > 0) args->channel->send_wait_us += ((uint64_t)(tv.tv_sec - args->channel->tv_send[buffSlot].tv_sec)*1000*1000 + tv.tv_usec - args->channel->tv_send[buffSlot].tv_usec);
if (args->opCount > 0) args->channel->send_byte += LOAD(sizesFifo+buffSlot);
args->channel->tv_send[buffSlot].tv_sec = tv.tv_sec;
args->channel->tv_send[buffSlot].tv_usec = tv.tv_usec;
#endif
STORE(sizesFifo+buffSlot, -1);
// Make sure size is reset to zero before we update the head.
__sync_synchronize();
@@ -318,6 +331,13 @@ ncclResult_t netSendProxy(struct ncclProxyArgs* args) {
int buffSlot = args->head%NCCL_STEPS;
NCCLCHECK(ncclNetTest(args->requests[buffSlot], &done, NULL));
if (done) {
#ifdef ENABLE_PROFILING
struct timeval tv;
gettimeofday(&tv, NULL);
if (args->opCount > 0) args->channel->send_us += ((uint64_t)(tv.tv_sec - args->channel->tv_send[buffSlot].tv_sec)*1000*1000 + tv.tv_usec - args->channel->tv_send[buffSlot].tv_usec);
args->channel->tv_send[buffSlot].tv_sec = 0;
args->channel->tv_send[buffSlot].tv_usec = 0;
#endif
args->head += args->sliceSteps;
STORE(&resources->sendMem->head, args->head);
args->idle = 0;
@@ -357,6 +377,9 @@ ncclResult_t netRecvProxy(struct ncclProxyArgs* args) {
if ((args->tail < args->head + NCCL_STEPS) && (args->tail < LOAD(sendHead) + NCCL_STEPS) && (args->tail < args->end)) {
int buffSlot = args->tail%NCCL_STEPS;
int sliceSize = stepSize * args->sliceSteps;
#ifdef ENABLE_PROFILING
gettimeofday(&args->channel->tv_recv[buffSlot], NULL);
#endif
NCCLCHECK(ncclNetIrecv(resources->netRecvComm, localBuff+buffSlot*stepSize, sliceSize, mhandle, args->requests+buffSlot));
if (args->requests[buffSlot] != NULL) {
args->tail += args->sliceSteps;
@@ -370,7 +393,19 @@ ncclResult_t netRecvProxy(struct ncclProxyArgs* args) {
if (done) {
args->head += args->sliceSteps;
if (args->protocol == NCCL_PROTO_SIMPLE) {
#ifdef ENABLE_PROFILING
struct timeval tv;
gettimeofday(&tv, NULL);
if (args->opCount > 0) args->channel->recv_byte += size;
if (args->opCount > 0) args->channel->recv_us += ((uint64_t)(tv.tv_sec - args->channel->tv_recv[buffSlot].tv_sec)*1000*1000 + tv.tv_usec - args->channel->tv_recv[buffSlot].tv_usec);
args->channel->tv_recv[buffSlot].tv_sec = tv.tv_sec;
args->channel->tv_recv[buffSlot].tv_usec = tv.tv_usec;
#endif
if (resources->useGdr) NCCLCHECK(ncclNetFlush(resources->netRecvComm, localBuff+buffSlot*stepSize, size, mhandle));
#ifdef ENABLE_PROFILING
gettimeofday(&tv, NULL);
if (args->opCount > 0) args->channel->recv_flush_us += ((uint64_t)(tv.tv_sec - args->channel->tv_recv[buffSlot].tv_sec)*1000*1000 + tv.tv_usec - args->channel->tv_recv[buffSlot].tv_usec);
#endif
STORE(&resources->recvMem->tail, args->head);
}
args->idle = 0;