Merge remote-tracking branch 'origin/develop' into 2.10.3
[ROCm/rccl commit: 8ee2b7932a]
This commit is contained in:
Wenkai Du
@@ -110,7 +110,7 @@ namespace {
|
||||
ACCUMULATE_COUNTER(directRecv);
|
||||
}
|
||||
#ifdef ENABLE_PROFILING
|
||||
if (tid == 0) __atomic_fetch_add(&(devProf->total_cycle), __builtin_amdgcn_s_memrealtime() - clk, __ATOMIC_SEQ_CST);
|
||||
if (tid == 0) devProf->elems[blockIdx.x].total_cycle += (__builtin_amdgcn_s_memrealtime() - clk);
|
||||
#endif
|
||||
}
|
||||
|
||||
|
||||
@@ -70,7 +70,7 @@ namespace {
|
||||
}
|
||||
}
|
||||
#ifdef ENABLE_PROFILING
|
||||
if (tid == 0) __atomic_fetch_add(&(devProf->total_cycle), __builtin_amdgcn_s_memrealtime() - clk, __ATOMIC_SEQ_CST);
|
||||
if (tid == 0) devProf->elems[blockIdx.x].total_cycle += (__builtin_amdgcn_s_memrealtime() - clk);
|
||||
#endif
|
||||
}
|
||||
}
|
||||
|
||||
@@ -15,12 +15,16 @@
|
||||
#define NCCL_SPINS_BEFORE_CHECK_ABORT 1000000
|
||||
|
||||
#define barrier_by_group() do { \
|
||||
const int w = threadIdx.x/WARP_SIZE; \
|
||||
const int wid = threadIdx.x%WARP_SIZE; \
|
||||
if (wid == 0) { \
|
||||
barrier_next[w] += nthreads/WARP_SIZE; \
|
||||
__atomic_fetch_add(barriers, 1, __ATOMIC_SEQ_CST); \
|
||||
while (LOAD(barriers) < barrier_next[w]) /* spin */; \
|
||||
if (nthreads == NCCL_MAX_NTHREADS) \
|
||||
__syncthreads(); \
|
||||
else { \
|
||||
const int w = threadIdx.x/WARP_SIZE; \
|
||||
const int wid = threadIdx.x%WARP_SIZE; \
|
||||
if (wid == 0) { \
|
||||
barrier_next[w] += nthreads/WARP_SIZE; \
|
||||
__atomic_fetch_add(barriers, 1, __ATOMIC_SEQ_CST); \
|
||||
while (LOAD(barriers) < barrier_next[w]) /* spin */; \
|
||||
} \
|
||||
} \
|
||||
} while (0)
|
||||
|
||||
@@ -153,12 +157,12 @@ struct PrimitivesWithoutDirect {
|
||||
|
||||
#ifdef ENABLE_PROFILING
|
||||
#define INIT_COUNTER \
|
||||
if (tid == 0) { t0 = __builtin_amdgcn_s_memrealtime(); ws = LOAD(&(devProf->wait_cycle[blockIdx.x])); }
|
||||
if (tid == 0) { t0 = __builtin_amdgcn_s_memrealtime(); ws = devProf->elems[blockIdx.x].wait_cycle; }
|
||||
|
||||
#define ACCUMULATE_COUNTER(prim) \
|
||||
if (tid == 0) { __atomic_fetch_add(&(devProf->prim##_cycle), __builtin_amdgcn_s_memrealtime() - t0 \
|
||||
+ ws - LOAD(&(devProf->wait_cycle[blockIdx.x])), __ATOMIC_SEQ_CST); \
|
||||
__atomic_fetch_add(&(devProf->prim##_byte), nelem * sizeof(T), __ATOMIC_SEQ_CST); }
|
||||
if (tid == 0) { 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); }
|
||||
#else
|
||||
#define INIT_COUNTER
|
||||
#define ACCUMULATE_COUNTER(prim)
|
||||
|
||||
@@ -293,11 +293,12 @@ class Primitives<T, RedOp, Fan, Direct, ProtoLL>:
|
||||
// Always waitSend in case of cleanup
|
||||
nelem = nelem < 0 ? 0 : nelem;
|
||||
#ifdef ENABLE_PROFILING
|
||||
uint64_t t0 = __builtin_amdgcn_s_memrealtime();
|
||||
uint64_t t0;
|
||||
if (tid == 0) t0 = __builtin_amdgcn_s_memrealtime();
|
||||
#endif
|
||||
if (SEND) waitSend(divUp(nelem, EltPerLine)*sizeof(ncclLLFifoLine));
|
||||
#ifdef ENABLE_PROFILING
|
||||
if (SEND && tid == 0) __atomic_fetch_add(&ncclShmem->comm.devProf->wait_cycle[blockIdx.x], __builtin_amdgcn_s_memrealtime() - t0, __ATOMIC_SEQ_CST);
|
||||
if (SEND && tid == 0) ncclShmem->comm.devProf->elems[blockIdx.x].wait_cycle = (__builtin_amdgcn_s_memrealtime() - t0);
|
||||
#endif
|
||||
|
||||
nelem -= tid*EltPerLine;
|
||||
|
||||
@@ -163,12 +163,13 @@ class Primitives<
|
||||
if (Dst && (flags & (DstBuf==Input ? RoleInput : RoleOutput)))
|
||||
ncclShmem->groups[group].dsts[0] = userBuff + dstIx + offset;
|
||||
#ifdef ENABLE_PROFILING
|
||||
uint64_t t0 = __builtin_amdgcn_s_memrealtime();
|
||||
uint64_t t0;
|
||||
if (tid == 0) t0 = __builtin_amdgcn_s_memrealtime();
|
||||
#endif
|
||||
waitPeer<DirectRecv, DirectSend, Recv, Send, Src, Dst>(dstIx, remoteOutIx, offset, sliceSize);
|
||||
subBarrier();
|
||||
#ifdef ENABLE_PROFILING
|
||||
if (tid == 0) __atomic_fetch_add(&ncclShmem->comm.devProf->wait_cycle[blockIdx.x], __builtin_amdgcn_s_memrealtime() - t0, __ATOMIC_SEQ_CST);
|
||||
if (tid == 0) 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
|
||||
|
||||
@@ -54,7 +54,7 @@ ncclResult_t parseList(const char* str, const char* elems[], int nelems, int* li
|
||||
|
||||
// Latencies in us, Bandwidths in GB/s
|
||||
// Tree { LL, LL128, Simple } , Ring { LL, LL128, Simple }
|
||||
static const float baseLat [NCCL_NUM_ALGORITHMS][NCCL_NUM_PROTOCOLS] = { { 39.0, 39.0, 49.0 }, { 21.0, 21.0, 30.0 }, { 37.9, 37.9, 40.4 } };
|
||||
static const float baseLat [NCCL_NUM_ALGORITHMS][NCCL_NUM_PROTOCOLS] = { { 12.0, 12.0, 17.0 }, { 12.0, 12.0, 17.0 }, { 12.0, 12.0, 17.0 } };
|
||||
|
||||
// NVLink, PCI, Network
|
||||
#define NCCL_HW_NVLINK 0
|
||||
@@ -63,11 +63,11 @@ static const float baseLat [NCCL_NUM_ALGORITHMS][NCCL_NUM_PROTOCOLS] = { { 39.0
|
||||
// Tree/Simple is the latency a 256kB chunk, which is ~ base lat + 256k/12GB/s (+ 256k/12GB/s for the network).
|
||||
static float hwLat [3][NCCL_NUM_ALGORITHMS][NCCL_NUM_PROTOCOLS] =
|
||||
{ /* NVLINK */
|
||||
{ /* Tree (LL/LL128/Simple)*/ { 2.5, 2.5, 5.5 }, /* Ring (LL/LL128/Simple)*/ { 2.5, 2.5, 5 }, /* CollNet (LL/LL128/Simple)*/ { 1.2, 1.2, 3.8 } },
|
||||
{ /* Tree (LL/LL128/Simple)*/ { 1.5, 1.5, 4.5 }, /* Ring (LL/LL128/Simple)*/ { 1.5, 1.5, 4.5 }, /* CollNet (LL/LL128/Simple)*/ { 1.5, 1.5, 4.5 } },
|
||||
/* PCI */
|
||||
{ /* Tree (LL/LL128/Simple)*/ { 2.2, 2.2, 5.7 }, /* Ring (LL/LL128/Simple)*/ { 1.3, 1.3, 1.9 }, /* CollNet (LL/LL128/Simple)*/ { 2.2, 2.2, 5.7 } },
|
||||
{ /* Tree (LL/LL128/Simple)*/ { 2.2, 2.2, 5.7 }, /* Ring (LL/LL128/Simple)*/ { 2.2, 2.2, 5.7 }, /* CollNet (LL/LL128/Simple)*/ { 2.2, 2.2, 5.7 } },
|
||||
/* NET */
|
||||
{ /* Tree (LL/LL128/Simple)*/ { 28.0, 28.0, 66.0 }, /* Ring (LL/LL128/Simple)*/ { 8.5, 8.5, 19.0 }, /* CollNet (LL/LL128/Simple)*/ { 9.8, 9.8, 19.5 } }
|
||||
{ /* Tree (LL/LL128/Simple)*/ { 40.0, 40.0, 50.0 }, /* Ring (LL/LL128/Simple)*/ { 4.0, 4.0, 25.0 }, /* CollNet (LL/LL128/Simple)*/ { 40.0, 40.0, 50.0 } }
|
||||
};
|
||||
|
||||
// LL128 max BW per channel
|
||||
@@ -75,7 +75,7 @@ static const double ll128MaxBwPerCh = 20.0;
|
||||
static const double llMaxBws[2][3] = { /* Volta-N1/Intel-N2/Intel-N4) */ {39.0, 39.0, 20.4}, /* Ampere-N1/AMD-N2/AMD-N4) */ {87.7, 22.5 /*avg of ring & tree*/, 19.0} };
|
||||
static const double perChMaxTreeBws[2][3] = { /* Volta (N1/N2/N4) */ {26.5, 18.5, 10.0}, /* Ampere (N1/N2/N4) */ {24.0, 23.6, 17.8} };
|
||||
|
||||
ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCompCap, struct ncclTopoGraph* treeGraph, struct ncclTopoGraph* ringGraph, struct ncclTopoGraph* collNetGraph) {
|
||||
ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCompCap, struct ncclTopoGraph* treeGraph, struct ncclTopoGraph* ringGraph, struct ncclTopoGraph* collNetGraph, int gcn) {
|
||||
int simpleDefaultThreads = (ringGraph->speedIntra*ringGraph->nChannels <= PCI_WIDTH) ? 256 : NCCL_SIMPLE_MAX_NTHREADS;
|
||||
comm->maxThreads[NCCL_ALGO_RING][NCCL_PROTO_SIMPLE] =
|
||||
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__)
|
||||
@@ -134,14 +134,10 @@ ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCom
|
||||
|
||||
// Various model refinements
|
||||
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__)
|
||||
if (a == NCCL_ALGO_RING && p == NCCL_PROTO_LL) busBw *= 1.0/5.0;
|
||||
if (a == NCCL_ALGO_RING && p == NCCL_PROTO_LL128) busBw = std::min(busBw * (ppn < 2 ? 0.7 : 0.92 /*120.0/128.0*/), ll128MaxBwPerCh*graphs[a]->nChannels);
|
||||
double maxTreeBw = comm->nNodes > 2 ?
|
||||
compCap80 && p == NCCL_PROTO_LL128 ? 105.0 : 80.0 :
|
||||
compCap80 && p == NCCL_PROTO_LL128 ? 130.0 : 110.0;
|
||||
if (a == NCCL_ALGO_TREE) busBw = std::min(busBw*.27, comm->nNodes > 1 ? 70.0 : 90.0);
|
||||
if (a == NCCL_ALGO_TREE && p == NCCL_PROTO_LL) busBw *= 1.0/2.3;
|
||||
if (a == NCCL_ALGO_TREE && p == NCCL_PROTO_LL128) busBw = std::min(busBw * (comm->nNodes == 1 ? 7.0/9.0 : 0.915 /*120.0/128.0*/), ll128MaxBwPerCh*graphs[a]->nChannels*7.0/9.0);
|
||||
if (a == NCCL_ALGO_RING && (p == NCCL_PROTO_LL || p == NCCL_PROTO_LL128)) busBw *= 0.05;
|
||||
if (a == NCCL_ALGO_TREE && p == NCCL_PROTO_SIMPLE) (nNodes == 2) ? busBw *= 0.33 : busBw *= 0.11;
|
||||
if (a == NCCL_ALGO_TREE && (p == NCCL_PROTO_LL || p == NCCL_PROTO_LL128)) busBw *= 0.04;
|
||||
if (gcn == 910 && p == NCCL_PROTO_LL && nNodes == 1 && nRanks == 16) busBw *= 5.9;
|
||||
#else
|
||||
if (compCap80) busBw = std::min(busBw, 235.0f);
|
||||
if (a == NCCL_ALGO_RING && p == NCCL_PROTO_LL) { busBw = std::min(llMaxBw, busBw * ((nNodes > 1 || coll == ncclFuncAllReduce || coll == ncclFuncReduce) ? 1.0/4.0 : 1.0/3.0)); }
|
||||
@@ -160,7 +156,7 @@ ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCom
|
||||
comm->latencies[coll][a][p] = baseLat[a][p];
|
||||
float intraLat = hwLat[intraHw[a]][a][p];
|
||||
float interLat = hwLat[NCCL_HW_NET][a][p];
|
||||
if (nNodes > 1 && p == NCCL_PROTO_LL) intraLat *= 1.8;
|
||||
//if (nNodes > 1 && p == NCCL_PROTO_LL) intraLat *= 1.8;
|
||||
if (a == NCCL_ALGO_RING) {
|
||||
float lat = hwLat[hw[a]][a][p];
|
||||
if ((coll == ncclFuncReduce || coll == ncclFuncBroadcast)) {
|
||||
@@ -286,15 +282,15 @@ ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCom
|
||||
// Trees are not perfectly sticking to the model for medium sizes. Applying a static correction
|
||||
// factor is not ideal but works quite well. Powers of two, 64 B to 128MB.
|
||||
static float treeCorrectionFactor[NCCL_NUM_PROTOCOLS][22] = {
|
||||
{ 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, .84, .49, .42, .60, .75, .87, .94, .94, .99, 1.0, 1.0 , 1.0 , 1.0 , 1.0 , 1.0 },
|
||||
{ 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, .84, .49, .42, .60, .75, .87, .94, .94, .99, 1.0, 1.0 , 1.0 , 1.0 , 1.0 , 1.0 },
|
||||
{ 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, .41, .27, .25, .39, .46, .72, .76, .87, .92, .97, 1.0, 1.0 , 1.0 , 1.0 , 1.0 , 1.0 }
|
||||
{ 0.7, 0.7, 0.7, 0.6, 0.6, 0.3, 0.9, 0.5, 0.5, 0.6, 0.5, 0.5, 0.8, 0.9, 0.8, 0.8, 0.8, 0.8, 0.8, 0.9, 0.9, 1.0, },
|
||||
{ 0.7, 0.7, 0.7, 0.6, 0.6, 0.3, 0.9, 0.5, 0.5, 0.6, 0.5, 0.5, 0.8, 0.9, 0.8, 0.8, 0.8, 0.8, 0.8, 0.9, 0.9, 1.0, },
|
||||
{ 0.4, 0.4, 0.3, 0.3, 0.2, 0.5, 0.5, 0.7, 0.2, 0.2, 0.3, 0.6, 0.7, 1.0, 1.3, 1.0, 1.2, 1.2, 1.1, 1.1, 1.2, 1.2, },
|
||||
};
|
||||
|
||||
static float ringCorrectionFactor[NCCL_NUM_PROTOCOLS][22] = {
|
||||
{ 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, .25, .41, .55, .56, .78, .94, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0 , 1.0 , 1.0 , 1.0 , 1.0 },
|
||||
{ 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, .25, .41, .55, .56, .78, .94, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0 , 1.0 , 1.0 , 1.0 , 1.0 },
|
||||
{ 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, .04, .08, .09, .09, .11, .13, .25, .40, .59, .76, .86, 1.0 , 1.0 , 1.0 , 1.0 , 1.0 }
|
||||
{ 0.4, 0.6, 0.6, 0.3, 0.2, 0.2, 0.2, 0.2, 0.4, 0.6, 0.7, 0.9, 1.4, 1.5, 1.0, 0.8, 0.7, 0.8, 0.8, 0.9, 0.9, 0.9, },
|
||||
{ 0.4, 0.6, 0.6, 0.3, 0.2, 0.2, 0.2, 0.2, 0.4, 0.6, 0.7, 0.9, 1.4, 1.5, 1.0, 0.8, 0.7, 0.8, 0.8, 0.9, 0.9, 0.9, },
|
||||
{ 0.6, 0.4, 0.4, 0.4, 0.2, 0.3, 0.2, 0.2, 0.2, 0.2, 0.2, 0.2, 0.2, 0.2, 0.2, 0.2, 0.2, 0.3, 0.4, 0.6, 0.8, 0.9, },
|
||||
};
|
||||
|
||||
ncclResult_t ncclTopoGetAlgoTime(struct ncclInfo* info, int algorithm, int protocol, int numPipeOps, float* time) {
|
||||
|
||||
@@ -18,8 +18,8 @@
|
||||
|
||||
// Convert volatile access to atomic
|
||||
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__)
|
||||
#define LOAD(VAR) __atomic_load_n((VAR), __ATOMIC_SEQ_CST)
|
||||
#define STORE(DST, SRC) __atomic_store_n((DST), (SRC), __ATOMIC_SEQ_CST)
|
||||
#define LOAD(VAR) __atomic_load_n((VAR), __ATOMIC_ACQUIRE)
|
||||
#define STORE(DST, SRC) __atomic_store_n((DST), (SRC), __ATOMIC_RELEASE)
|
||||
#else
|
||||
#define LOAD(VAR) *(VAR)
|
||||
#define STORE(DST, SRC) *(DST) = (SRC)
|
||||
@@ -283,11 +283,11 @@ static_assert(sizeof(struct ncclChannel) == 0x80*sizeof(int), "ncclChannel must
|
||||
#pragma pack(pop) /* restore original alignment from stack */
|
||||
|
||||
#ifdef ENABLE_PROFILING
|
||||
struct ncclProf {
|
||||
struct ncclProfElem {
|
||||
union {
|
||||
struct {
|
||||
uint64_t total_cycle;
|
||||
uint64_t wait_cycle[MAXCHANNELS]; // total wait cycle
|
||||
uint64_t wait_cycle; // total wait cycle
|
||||
// primtive cycles
|
||||
uint64_t send_cycle;
|
||||
uint64_t directSend_cycle;
|
||||
@@ -318,6 +318,10 @@ struct ncclProf {
|
||||
int data[0x80];
|
||||
};
|
||||
};
|
||||
|
||||
struct ncclProf {
|
||||
struct ncclProfElem elems[MAXCHANNELS];
|
||||
};
|
||||
#endif
|
||||
|
||||
#ifdef ENABLE_COLLTRACE
|
||||
|
||||
@@ -106,7 +106,7 @@ ncclResult_t ncclTopoPreset(struct ncclComm* comm,
|
||||
ncclResult_t ncclTopoPostset(struct ncclComm* comm, int* firstRanks, int* treePatterns,
|
||||
struct ncclTopoRanks** allTopoRanks, int* rings, struct ncclTopoGraph* collNetGraph, int nc);
|
||||
|
||||
ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCompCap, struct ncclTopoGraph* treeGraph, struct ncclTopoGraph* ringGraph, struct ncclTopoGraph* collNetGraph);
|
||||
ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCompCap, struct ncclTopoGraph* treeGraph, struct ncclTopoGraph* ringGraph, struct ncclTopoGraph* collNetGraph, int gcn);
|
||||
#include "info.h"
|
||||
ncclResult_t ncclTopoGetAlgoTime(struct ncclInfo* info, int algorithm, int protocol, int numPipeOps, float* time);
|
||||
|
||||
|
||||
+42
-12
@@ -312,9 +312,39 @@ 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_cycle = 0;
|
||||
uint64_t total_cycle = 0, wait_cycle = 0, send_cycle = 0, directSend_cycle = 0, recv_cycle = 0, \
|
||||
directRecv_cycle = 0, copySend_cycle = 0, directCopySend_cycle = 0, recvCopySend_cycle = 0, directRecvCopySend_cycle = 0, \
|
||||
recvReduceCopy_cycle = 0, recvReduceSend_cycle = 0, recvReduceCopySend_cycle = 0, directRecvReduceCopySend_cycle = 0, \
|
||||
send_byte = 0, directSend_byte = 0, recv_byte = 0, directRecv_byte = 0, copySend_byte = 0, directCopySend_byte = 0, \
|
||||
recvCopySend_byte = 0, directRecvCopySend_byte = 0, recvReduceCopy_byte = 0, recvReduceSend_byte = 0, \
|
||||
recvReduceCopySend_byte = 0, directRecvReduceCopySend_byte = 0;
|
||||
for (int chan=0; chan<comm->nChannels; chan++) {
|
||||
wait_cycle += prof->wait_cycle[chan];
|
||||
total_cycle += prof->elems[chan].total_cycle;
|
||||
wait_cycle += prof->elems[chan].wait_cycle;
|
||||
send_cycle += prof->elems[chan].send_cycle;
|
||||
directSend_cycle += prof->elems[chan].directSend_cycle;
|
||||
recv_cycle += prof->elems[chan].recv_cycle;
|
||||
directRecv_cycle += prof->elems[chan].directRecv_cycle;
|
||||
copySend_cycle += prof->elems[chan].copySend_cycle;
|
||||
directCopySend_cycle += prof->elems[chan].directCopySend_cycle;
|
||||
recvCopySend_cycle += prof->elems[chan].recvCopySend_cycle;
|
||||
directRecvCopySend_cycle += prof->elems[chan].directRecvCopySend_cycle;
|
||||
recvReduceCopy_cycle += prof->elems[chan].recvReduceCopy_cycle;
|
||||
recvReduceSend_cycle += prof->elems[chan].recvReduceSend_cycle;
|
||||
recvReduceCopySend_cycle += prof->elems[chan].recvReduceCopySend_cycle;
|
||||
directRecvReduceCopySend_cycle += prof->elems[chan].directRecvReduceCopySend_cycle;
|
||||
send_byte += prof->elems[chan].send_byte;
|
||||
directSend_byte += prof->elems[chan].directSend_byte;
|
||||
recv_byte += prof->elems[chan].recv_byte;
|
||||
directRecv_byte += prof->elems[chan].directRecv_byte;
|
||||
copySend_byte += prof->elems[chan].copySend_byte;
|
||||
directCopySend_byte += prof->elems[chan].directCopySend_byte;
|
||||
recvCopySend_byte += prof->elems[chan].recvCopySend_byte;
|
||||
directRecvCopySend_byte += prof->elems[chan].directRecvCopySend_byte;
|
||||
recvReduceCopy_byte += prof->elems[chan].recvReduceCopy_byte;
|
||||
recvReduceSend_byte += prof->elems[chan].recvReduceSend_byte;
|
||||
recvReduceCopySend_byte += prof->elems[chan].recvReduceCopySend_byte;
|
||||
directRecvReduceCopySend_byte += prof->elems[chan].directRecvReduceCopySend_byte;
|
||||
}
|
||||
#define VEGA_GPU_RTC_FREQUENCY 2.5E7
|
||||
if (comm->rank == 0) {
|
||||
@@ -322,16 +352,16 @@ static ncclResult_t commFree(ncclComm_t comm) {
|
||||
INFO(NCCL_INIT, "# %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)");
|
||||
}
|
||||
INFO(NCCL_INIT, "# %4d %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,
|
||||
comm->rank, (double)total_cycle/VEGA_GPU_RTC_FREQUENCY/comm->nChannels,
|
||||
(double)wait_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,
|
||||
(prof->directRecvReduceCopySend_cycle) ? (double)prof->directRecvReduceCopySend_byte*comm->nChannels/((double)prof->directRecvReduceCopySend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0,
|
||||
(prof->directRecvCopySend_cycle) ? (double)prof->directRecvCopySend_byte*comm->nChannels/((double)prof->directRecvCopySend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0,
|
||||
(prof->directRecv_cycle) ? (double)prof->directRecv_byte*comm->nChannels/((double)prof->directRecv_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0,
|
||||
(prof->copySend_cycle) ? (double)prof->copySend_byte*comm->nChannels/((double)prof->copySend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0,
|
||||
(prof->recv_cycle) ? (double)prof->recv_byte*comm->nChannels/((double)prof->recv_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0,
|
||||
(prof->recvCopySend_cycle) ? (double)prof->recvCopySend_byte*comm->nChannels/((double)prof->recvCopySend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0);
|
||||
(send_cycle) ? (double)send_byte*comm->nChannels/((double)send_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0,
|
||||
(recvReduceSend_cycle) ? (double)recvReduceSend_byte*comm->nChannels/((double)recvReduceSend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0,
|
||||
(directRecvReduceCopySend_cycle) ? (double)directRecvReduceCopySend_byte*comm->nChannels/((double)directRecvReduceCopySend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0,
|
||||
(directRecvCopySend_cycle) ? (double)directRecvCopySend_byte*comm->nChannels/((double)directRecvCopySend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0,
|
||||
(directRecv_cycle) ? (double)directRecv_byte*comm->nChannels/((double)directRecv_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0,
|
||||
(copySend_cycle) ? (double)copySend_byte*comm->nChannels/((double)copySend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0,
|
||||
(recv_cycle) ? (double)recv_byte*comm->nChannels/((double)recv_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0,
|
||||
(recvCopySend_cycle) ? (double)recvCopySend_byte*comm->nChannels/((double)recvCopySend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0);
|
||||
free(prof);
|
||||
CUDACHECK(hipFree(comm->hostDevComm.devProf));
|
||||
|
||||
@@ -1114,7 +1144,7 @@ collnet_cleanup:
|
||||
TRACE(NCCL_INIT, "rank %d nranks %d - CONNECTED %d RINGS AND TREES", rank, nranks, comm->nChannels);
|
||||
|
||||
// Compute time models for algorithm and protocol combinations
|
||||
NCCLCHECK(ncclTopoTuneModel(comm, minCompCap, maxCompCap, &treeGraph, &ringGraph, &collNetGraph));
|
||||
NCCLCHECK(ncclTopoTuneModel(comm, minCompCap, maxCompCap, &treeGraph, &ringGraph, &collNetGraph, comm->topo->nodes[GPU].nodes[0].gpu.gcn));
|
||||
|
||||
// Compute nChannels per peer for p2p
|
||||
NCCLCHECK(ncclTopoComputeP2pChannels(comm));
|
||||
|
||||
@@ -0,0 +1,143 @@
|
||||
/*
|
||||
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
// Helper macro for checking HSA calls
|
||||
#define HSA_CHECK(cmd) \
|
||||
do { \
|
||||
hsa_status_t error = (cmd); \
|
||||
if (error != HSA_STATUS_SUCCESS) { \
|
||||
const char* errString = NULL; \
|
||||
hsa_status_string(error, &errString); \
|
||||
std::cerr << "Encountered HSA error (" << errString << ") at line " \
|
||||
<< __LINE__ << " in file " << __FILE__ << "\n"; \
|
||||
exit(-1); \
|
||||
} \
|
||||
} while (0)
|
||||
|
||||
// Structure to hold HSA agent information
|
||||
struct AgentData
|
||||
{
|
||||
bool isInitialized;
|
||||
std::vector<hsa_agent_t> cpuAgents;
|
||||
std::vector<hsa_agent_t> gpuAgents;
|
||||
std::vector<int> closestNumaNode;
|
||||
};
|
||||
|
||||
// Simple callback function to return any memory pool for an agent
|
||||
hsa_status_t MemPoolInfoCallback(hsa_amd_memory_pool_t pool, void *data)
|
||||
{
|
||||
hsa_amd_memory_pool_t* poolData = reinterpret_cast<hsa_amd_memory_pool_t*>(data);
|
||||
|
||||
// Check memory pool flags
|
||||
uint32_t poolFlags;
|
||||
HSA_CHECK(hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &poolFlags));
|
||||
|
||||
// Only consider coarse-grained pools
|
||||
if (!(poolFlags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINED)) return HSA_STATUS_SUCCESS;
|
||||
|
||||
*poolData = pool;
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
// Callback function to gather HSA agent information
|
||||
hsa_status_t AgentInfoCallback(hsa_agent_t agent, void* data)
|
||||
{
|
||||
AgentData* agentData = reinterpret_cast<AgentData*>(data);
|
||||
|
||||
// Get the device type
|
||||
hsa_device_type_t deviceType;
|
||||
HSA_CHECK(hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &deviceType));
|
||||
if (deviceType == HSA_DEVICE_TYPE_CPU)
|
||||
agentData->cpuAgents.push_back(agent);
|
||||
if (deviceType == HSA_DEVICE_TYPE_GPU)
|
||||
{
|
||||
agentData->gpuAgents.push_back(agent);
|
||||
agentData->closestNumaNode.push_back(0);
|
||||
}
|
||||
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
AgentData& GetAgentData()
|
||||
{
|
||||
static AgentData agentData = {};
|
||||
|
||||
if (!agentData.isInitialized)
|
||||
{
|
||||
agentData.isInitialized = true;
|
||||
|
||||
// Add all detected agents to the list
|
||||
HSA_CHECK(hsa_iterate_agents(AgentInfoCallback, &agentData));
|
||||
|
||||
// Loop over each GPU
|
||||
for (uint32_t i = 0; i < agentData.gpuAgents.size(); i++)
|
||||
{
|
||||
// Collect memory pool
|
||||
hsa_amd_memory_pool_t pool;
|
||||
HSA_CHECK(hsa_amd_agent_iterate_memory_pools(agentData.gpuAgents[i], MemPoolInfoCallback, &pool));
|
||||
|
||||
// Loop over each CPU agent and check distance
|
||||
int bestDistance = -1;
|
||||
for (uint32_t j = 0; j < agentData.cpuAgents.size(); j++)
|
||||
{
|
||||
// Determine number of hops from GPU memory pool to CPU agent
|
||||
uint32_t hops = 0;
|
||||
HSA_CHECK(hsa_amd_agent_memory_pool_get_info(agentData.cpuAgents[j],
|
||||
pool,
|
||||
HSA_AMD_AGENT_MEMORY_POOL_INFO_NUM_LINK_HOPS,
|
||||
&hops));
|
||||
// Gather link info
|
||||
hsa_amd_memory_pool_link_info_t* link_info =
|
||||
(hsa_amd_memory_pool_link_info_t *)malloc(hops * sizeof(hsa_amd_memory_pool_link_info_t));
|
||||
HSA_CHECK(hsa_amd_agent_memory_pool_get_info(agentData.cpuAgents[j],
|
||||
pool,
|
||||
HSA_AMD_AGENT_MEMORY_POOL_INFO_LINK_INFO,
|
||||
link_info));
|
||||
int numaDist = 0;
|
||||
for (int k = 0; k < hops; k++)
|
||||
{
|
||||
numaDist += link_info[k].numa_distance;
|
||||
}
|
||||
if (bestDistance == -1 || numaDist < bestDistance)
|
||||
{
|
||||
agentData.closestNumaNode[i] = j;
|
||||
bestDistance = numaDist;
|
||||
}
|
||||
free(link_info);
|
||||
}
|
||||
}
|
||||
}
|
||||
return agentData;
|
||||
}
|
||||
|
||||
// Returns closest CPU NUMA node to provided GPU
|
||||
// NOTE: This assumes HSA GPU indexing is similar to HIP GPU indexing
|
||||
int GetClosestNumaNode(int gpuIdx)
|
||||
{
|
||||
AgentData& agentData = GetAgentData();
|
||||
if (gpuIdx < 0 || gpuIdx >= agentData.closestNumaNode.size())
|
||||
{
|
||||
printf("[ERROR] GPU index out is out of bounds\n");
|
||||
exit(1);
|
||||
}
|
||||
return agentData.closestNumaNode[gpuIdx];
|
||||
}
|
||||
@@ -6,7 +6,7 @@ endif
|
||||
HIPCC=$(HIP_PATH)/bin/hipcc
|
||||
|
||||
EXE=TransferBench
|
||||
CXXFLAGS = -O3 -I../../src/include -I. -lnuma
|
||||
CXXFLAGS = -O3 -I../../src/include -I. -lnuma -L$(HIP_PATH)/../hsa/lib -lhsa-runtime64
|
||||
|
||||
all: $(EXE)
|
||||
|
||||
|
||||
@@ -24,6 +24,7 @@ THE SOFTWARE.
|
||||
// on the same node
|
||||
|
||||
#include "TransferBench.hpp"
|
||||
#include "GetClosestNumaNode.hpp"
|
||||
#include <numa.h>
|
||||
#include <numaif.h>
|
||||
#include <stack>
|
||||
@@ -270,10 +271,12 @@ int main(int argc, char **argv)
|
||||
// Report timings
|
||||
totalCpuTime = totalCpuTime / (1.0 * ev.numIterations) * 1000;
|
||||
double totalBandwidthGbs = (numLinks * N * sizeof(float) / 1.0E6) / totalCpuTime;
|
||||
double maxGpuTime = 0;
|
||||
for (int i = 0; i < numLinks; i++)
|
||||
{
|
||||
double linkDurationMsec = links[i].totalTime / (1.0 * ev.numIterations);
|
||||
double linkBandwidthGbs = (N * sizeof(float) / 1.0E9) / linkDurationMsec * 1000.0f;
|
||||
maxGpuTime = std::max(maxGpuTime, linkDurationMsec);
|
||||
if (!ev.outputToCsv)
|
||||
{
|
||||
printf(" Link %02d: %c%02d -> [%cPU %02d:%02d] -> %c%02d | %9.3f GB/s | %8.3f ms | %-16s",
|
||||
@@ -310,7 +313,8 @@ int main(int argc, char **argv)
|
||||
// Display aggregate statistics
|
||||
if (!ev.outputToCsv)
|
||||
{
|
||||
printf(" Aggregate Bandwidth (CPU timed) | %9.3f GB/s | %8.3f ms |\n", totalBandwidthGbs, totalCpuTime);
|
||||
printf(" Aggregate Bandwidth (CPU timed) | %9.3f GB/s | %8.3f ms | Overhead: %.3f ms\n", totalBandwidthGbs, totalCpuTime,
|
||||
totalCpuTime - maxGpuTime);
|
||||
}
|
||||
else
|
||||
{
|
||||
@@ -514,7 +518,7 @@ void GenerateConfigFile(char const* cfgFile, int numBlocks)
|
||||
fprintf(fp, "# GPU 0 Gather\n");
|
||||
fprintf(fp, "%d %d", numGpuDevices-1, numBlocks);
|
||||
for (int i = 1; i < numGpuDevices; i++)
|
||||
fprintf(fp, " (G%d->G%d->G%d)", 0, i, 0);
|
||||
fprintf(fp, " (G%d->G%d->G%d)", i, 0, 0);
|
||||
fprintf(fp, "\n\n");
|
||||
|
||||
// Full stress test
|
||||
@@ -533,17 +537,16 @@ void GenerateConfigFile(char const* cfgFile, int numBlocks)
|
||||
|
||||
void DisplayTopology()
|
||||
{
|
||||
printf("\nDetected topology:\n");
|
||||
int numGpuDevices;
|
||||
HIP_CALL(hipGetDeviceCount(&numGpuDevices));
|
||||
|
||||
printf("\nDetected topology: %d CPU NUMA node(s) %d GPU device(s)\n", numa_num_configured_nodes(), numGpuDevices);
|
||||
printf(" |");
|
||||
for (int j = 0; j < numGpuDevices; j++)
|
||||
printf(" GPU %02d |", j);
|
||||
printf(" PCIe Bus ID\n");
|
||||
printf(" PCIe Bus ID | Closest NUMA\n");
|
||||
for (int j = 0; j <= numGpuDevices; j++)
|
||||
printf("--------+");
|
||||
printf("-------------\n");
|
||||
printf("--------------+-------------\n");
|
||||
|
||||
char pciBusId[20];
|
||||
for (int i = 0; i < numGpuDevices; i++)
|
||||
@@ -567,7 +570,7 @@ void DisplayTopology()
|
||||
}
|
||||
}
|
||||
HIP_CALL(hipDeviceGetPCIBusId(pciBusId, 20, i));
|
||||
printf(" %s\n", pciBusId);
|
||||
printf(" %11s | %d \n", pciBusId, GetClosestNumaNode(i));
|
||||
}
|
||||
}
|
||||
|
||||
@@ -670,7 +673,7 @@ void ParseLinks(char* line, int numCpus, int numGpus, std::vector<Link>& links)
|
||||
// Method 1: Take in triples (srcMem, exeMem, dstMem)
|
||||
int numBlocksToUse;
|
||||
iss >> numBlocksToUse;
|
||||
if (numBlocksToUse <= 0)
|
||||
if (numBlocksToUse <= 0 || iss.fail())
|
||||
{
|
||||
printf("Parsing error: Number of blocks to use (%d) must be greater than 0\n", numBlocksToUse);
|
||||
exit(1);
|
||||
@@ -679,6 +682,11 @@ void ParseLinks(char* line, int numCpus, int numGpus, std::vector<Link>& links)
|
||||
for (int i = 0; i < numLinks; i++)
|
||||
{
|
||||
iss >> srcMem >> exeMem >> dstMem;
|
||||
if (iss.fail())
|
||||
{
|
||||
printf("Parsing error: Unable to read valid Link triplet (possibly missing a SRC or EXE or DST)\n");
|
||||
exit(1);
|
||||
}
|
||||
ParseMemType(srcMem, numCpus, numGpus, &links[i].srcMemType, &links[i].srcIndex);
|
||||
ParseMemType(exeMem, numCpus, numGpus, &links[i].exeMemType, &links[i].exeIndex);
|
||||
ParseMemType(dstMem, numCpus, numGpus, &links[i].dstMemType, &links[i].dstIndex);
|
||||
@@ -699,6 +707,11 @@ void ParseLinks(char* line, int numCpus, int numGpus, std::vector<Link>& links)
|
||||
for (int i = 0; i < numLinks; i++)
|
||||
{
|
||||
iss >> srcMem >> exeMem >> dstMem >> links[i].numBlocksToUse;
|
||||
if (iss.fail())
|
||||
{
|
||||
printf("Parsing error: Unable to read valid Link quadruple (possibly missing a SRC or EXE or DST or #CU)\n");
|
||||
exit(1);
|
||||
}
|
||||
ParseMemType(srcMem, numCpus, numGpus, &links[i].srcMemType, &links[i].srcIndex);
|
||||
ParseMemType(exeMem, numCpus, numGpus, &links[i].exeMemType, &links[i].exeIndex);
|
||||
ParseMemType(dstMem, numCpus, numGpus, &links[i].dstMemType, &links[i].dstIndex);
|
||||
|
||||
@@ -264,4 +264,4 @@ __device__ __forceinline__ void ReduceOrCopyMulti(const int tid, const int nthre
|
||||
ReduceCopyMulti<FUNC, T, 1, MINSRCS, MAXSRCS, MINDSTS, MAXDSTS>(w, nw, t, nsrcs, srcs, ndsts, dsts, offset, Nrem);
|
||||
}
|
||||
|
||||
#endif // COMMON_KERNEL_H_
|
||||
#endif
|
||||
@@ -33,16 +33,16 @@ THE SOFTWARE.
|
||||
#include <hip/hip_runtime.h>
|
||||
#include "copy_kernel.h"
|
||||
|
||||
#define MAX_GPU 8
|
||||
#define MAX_GPU 16
|
||||
#define MAX_WORKGROUPS 32
|
||||
#define THREADS 256
|
||||
#define NGPUS 2
|
||||
|
||||
#define COPY_UNROLL 4
|
||||
#define REDUCE_UNROLL 2
|
||||
#define DOUBLECOPY_UNROLL 2
|
||||
#define DOUBLECOPYLOCAL_UNROLL 2
|
||||
#define REDUCECOPY_UNROLL 2
|
||||
#define ALL2ALL_UNROLL 2
|
||||
|
||||
|
||||
|
||||
@@ -59,18 +59,15 @@ THE SOFTWARE.
|
||||
#define RTC_CLOCK_FREQ_DEFAULT 2.7E7
|
||||
|
||||
struct transfer_data_t {
|
||||
// Buffers for all OPs except all to all
|
||||
float *dest0[MAX_WORKGROUPS]; //remote fine grain
|
||||
float *src0[MAX_WORKGROUPS]; //local fine grain
|
||||
float *dest1[MAX_WORKGROUPS]; //local coarse grain
|
||||
float *dest2[MAX_WORKGROUPS]; //local fine grain
|
||||
float *src1[MAX_WORKGROUPS]; //local coarse grain
|
||||
// Buffers for all to all
|
||||
const float *srcs[MAX_WORKGROUPS][MAX_GPU];
|
||||
float *dsts[MAX_WORKGROUPS][MAX_GPU];
|
||||
int N;
|
||||
int gpu;
|
||||
int ngpu;
|
||||
uint64_t *remOpCount;
|
||||
};
|
||||
|
||||
struct profiling_data_t {
|
||||
@@ -100,11 +97,10 @@ enum Ops {
|
||||
OP_REDUCE,
|
||||
OP_REDUCECOPY,
|
||||
OP_READ,
|
||||
OP_ALL2ALL,
|
||||
NUM_OPS,
|
||||
};
|
||||
|
||||
template<int op, int NGPUS>
|
||||
template<int op, int sync>
|
||||
__global__ void flag_sync_kernel(struct transfer_data_t* transfer_data, struct profiling_data_t* profiling_data, uint64_t opCount) {
|
||||
size_t tid = threadIdx.x;
|
||||
uint64_t curr_time;
|
||||
@@ -114,9 +110,19 @@ __global__ void flag_sync_kernel(struct transfer_data_t* transfer_data, struct p
|
||||
const float *srcs[NGPUS];
|
||||
float *dsts[NGPUS];
|
||||
|
||||
// signal self ready and wait until all GPUs are ready
|
||||
if (tid == 0) {
|
||||
curr_time = __builtin_amdgcn_s_memrealtime();
|
||||
__atomic_fetch_add(&transfer_data->remOpCount[transfer_data->gpu], 1, __ATOMIC_SEQ_CST);
|
||||
if (sync) {
|
||||
for (int i = 0; i < transfer_data->ngpu; i++) {
|
||||
while (LOAD(&transfer_data->remOpCount[i]) < opCount) {};
|
||||
}
|
||||
}
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
if (tid == 0)
|
||||
curr_time = __builtin_amdgcn_s_memrealtime();
|
||||
|
||||
if (op == OP_COPY) {
|
||||
srcs[0] = transfer_data->src0[bid];
|
||||
@@ -167,36 +173,32 @@ __global__ void flag_sync_kernel(struct transfer_data_t* transfer_data, struct p
|
||||
ReduceOrCopyMulti<COPY_UNROLL, FuncPassA<float>, float, 1, 1, 1, 1>(threadIdx.x, THREADS,
|
||||
1, srcs, 1, dsts, n);
|
||||
}
|
||||
if (op == OP_ALL2ALL) {
|
||||
for (int i = 0; i < NGPUS; i++) {
|
||||
srcs[i] = transfer_data->srcs[bid][i];
|
||||
dsts[i] = transfer_data->dsts[bid][i];
|
||||
}
|
||||
ReduceOrCopyMulti<ALL2ALL_UNROLL, FuncSum<float>, float, 1, NGPUS, 1, NGPUS>(tid, THREADS,
|
||||
NGPUS, srcs, NGPUS, dsts, n);
|
||||
}
|
||||
|
||||
__syncthreads();
|
||||
|
||||
if (tid == 0) {
|
||||
__atomic_fetch_add(&(profiling_data->write_cycles[bid]), __builtin_amdgcn_s_memrealtime() - curr_time, __ATOMIC_SEQ_CST);
|
||||
// for all to all, read and write n itmes to all other GPUs, thus "n * sizeof(float) * (transfer_data->ngpu - 1) * 2" bytes
|
||||
if (op == OP_ALL2ALL) __atomic_fetch_add(&(profiling_data->bytes_transferred[bid]), n * sizeof(float) * (transfer_data->ngpu - 1) * 2, __ATOMIC_SEQ_CST);
|
||||
else __atomic_fetch_add(&(profiling_data->bytes_transferred[bid]), n * sizeof(float), __ATOMIC_SEQ_CST);
|
||||
__atomic_fetch_add(&(profiling_data->bytes_transferred[bid]), n * sizeof(float), __ATOMIC_SEQ_CST);
|
||||
}
|
||||
}
|
||||
|
||||
typedef void(*flag_sync_kernel_t)(struct transfer_data_t* transfer_data, struct profiling_data_t* profiling_data, uint64_t opCount);
|
||||
|
||||
static flag_sync_kernel_t const flagSyncKerns[NUM_OPS+1] = {
|
||||
flag_sync_kernel<OP_COPY, 2>,
|
||||
flag_sync_kernel<OP_LOCALCOPY, 2>,
|
||||
flag_sync_kernel<OP_DOUBLECOPY, 2>,
|
||||
flag_sync_kernel<OP_DOUBLECOPYLOCAL, 2>,
|
||||
flag_sync_kernel<OP_REDUCE, 2>,
|
||||
flag_sync_kernel<OP_REDUCECOPY, 2>,
|
||||
flag_sync_kernel<OP_READ, 2>,
|
||||
flag_sync_kernel<OP_ALL2ALL, 4>,
|
||||
flag_sync_kernel<OP_ALL2ALL, 8>,
|
||||
static flag_sync_kernel_t const flagSyncKerns[NUM_OPS*2] = {
|
||||
flag_sync_kernel<OP_COPY, 0>,
|
||||
flag_sync_kernel<OP_COPY, 1>,
|
||||
flag_sync_kernel<OP_LOCALCOPY, 0>,
|
||||
flag_sync_kernel<OP_LOCALCOPY, 1>,
|
||||
flag_sync_kernel<OP_DOUBLECOPY, 0>,
|
||||
flag_sync_kernel<OP_DOUBLECOPY, 1>,
|
||||
flag_sync_kernel<OP_DOUBLECOPYLOCAL, 0>,
|
||||
flag_sync_kernel<OP_DOUBLECOPYLOCAL, 1>,
|
||||
flag_sync_kernel<OP_REDUCE, 0>,
|
||||
flag_sync_kernel<OP_REDUCE, 1>,
|
||||
flag_sync_kernel<OP_REDUCECOPY, 0>,
|
||||
flag_sync_kernel<OP_REDUCECOPY, 1>,
|
||||
flag_sync_kernel<OP_READ, 0>,
|
||||
flag_sync_kernel<OP_READ, 1>,
|
||||
};
|
||||
|
||||
__global__ void initTestDataKernel(float* data, const size_t N, const int gpu) {
|
||||
@@ -230,16 +232,18 @@ static void setupPeers(uint32_t *info, bool* is_xgmi) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
for (int j = 0; j < deviceCnt; j++) {
|
||||
if (i != j) {
|
||||
int p2p;
|
||||
int p2p;
|
||||
HIPCHECK(hipDeviceCanAccessPeer(&p2p, i, j));
|
||||
if (!p2p) {
|
||||
printf("Cannot enable peer access between device %d and %d. You may use HIP_VISIBLE_DEVICES to limit GPUs.\n",
|
||||
i, j);
|
||||
i, j);
|
||||
exit(-1);
|
||||
}
|
||||
HIPCHECK(hipDeviceEnablePeerAccess(j, 0));
|
||||
uint32_t linktype;
|
||||
HIPCHECK(hipExtGetLinkTypeAndHopCount(i, j, &linktype, &info[i*deviceCnt+j]));
|
||||
hipError_t error = hipExtGetLinkTypeAndHopCount(i, j, &linktype, &info[i*deviceCnt+j]);
|
||||
if (error != hipSuccess)
|
||||
*is_xgmi = 0;
|
||||
if (linktype != 4 || info[i*deviceCnt+j] != 1) *is_xgmi = 0;
|
||||
}
|
||||
else
|
||||
@@ -266,7 +270,9 @@ static void parseChordalRing(char **str) {
|
||||
int count = 0;
|
||||
for (int n = 0; n<ngpus; n++) {
|
||||
uint32_t linktype, hop;
|
||||
HIPCHECK(hipExtGetLinkTypeAndHopCount(i, n, &linktype, &hop));
|
||||
hipError_t error = hipExtGetLinkTypeAndHopCount(i, n, &linktype, &hop);
|
||||
if (error != hipSuccess)
|
||||
return;
|
||||
if (linktype != 4 || hop != 1) continue;
|
||||
sum -= n;
|
||||
count ++;
|
||||
@@ -368,11 +374,11 @@ static const char* link_type_name[] = {"HT", "QPI", "PCIE", "IB", "XGMI"};
|
||||
int main(int argc,char* argv[])
|
||||
{
|
||||
if (cmdOptionExists(argv, argv + argc, "-h")) {
|
||||
printf("./rccl_prim_test -w num_workgroups -p copy|localcopy|doublecopy|doublecopylocal|reduce|reducecopy|all2all -i iterations -n bytes -r \"0 1 2 3|3 2 1 0\"\n");
|
||||
printf("./rccl_prim_test -w num_workgroups -p copy|localcopy|doublecopy|doublecopylocal|reduce|reducecopy|all -i iterations -n bytes -r \"0 1 2 3|3 2 1 0\"\n");
|
||||
exit(0);
|
||||
}
|
||||
|
||||
int workgroups = 1;
|
||||
int workgroups = 0;
|
||||
char *wg = getCmdOption(argv, argv + argc, "-w");
|
||||
if (wg)
|
||||
workgroups = atol(wg);
|
||||
@@ -391,10 +397,16 @@ int main(int argc,char* argv[])
|
||||
printf("Benchmarking using %ld bytes\n", nBytes);
|
||||
uint64_t N = nBytes/sizeof(float);
|
||||
|
||||
int sync = 1;
|
||||
char *s = getCmdOption(argv, argv + argc, "-s");
|
||||
if (s)
|
||||
sync = atol(s);
|
||||
if (sync) printf("Sync all GPUs before operation\n");
|
||||
|
||||
char *r = getCmdOption(argv, argv + argc, "-r");
|
||||
if (r) printf("User specified ring topology: %s\n", r);
|
||||
|
||||
const char *ops[] = {"copy", "localcopy", "doublecopy", "doublecopylocal", "reduce", "reducecopy", "read", "all2all"};
|
||||
const char *ops[] = {"copy", "localcopy", "doublecopy", "doublecopylocal", "reduce", "reducecopy", "read", "all"};
|
||||
char *prim = getCmdOption(argv, argv + argc, "-p");
|
||||
int op = NUM_OPS, begin_op, end_op;
|
||||
if (prim) {
|
||||
@@ -419,11 +431,24 @@ int main(int argc,char* argv[])
|
||||
bool is_xgmi;
|
||||
char *cr8g = 0;
|
||||
static const char *ring_4p3l = "0 1 2 3|0 1 3 2|0 2 1 3|0 2 3 1|0 3 1 2|0 3 2 1";
|
||||
static const char *ring_8p1h = "0 1 3 2 4 5 7 6|6 7 5 4 2 3 1 0|0 1 5 4 6 7 3 2|2 3 7 6 4 5 1 0";
|
||||
static const char *ring_16p1h = "0 1 3 2 6 7 15 14 10 11 9 8 12 13 5 4|0 1 2 3 7 6 13 12 8 9 10 11 15 14 5 4|0 2 3 7 6 14 15 11 10 8 9 13 12 4 5 1|4 5 13 12 8 9 11 10 14 15 7 6 2 3 1 0|4 5 14 15 11 10 9 8 12 13 6 7 3 2 1 0|1 5 4 12 13 9 8 10 11 15 14 6 7 3 2 0";
|
||||
setupPeers(connection_info, &is_xgmi);
|
||||
parseChordalRing(&cr8g);
|
||||
if (nGpu == 4 && is_xgmi) r = (char *)ring_4p3l;
|
||||
if (nGpu == 8 && cr8g) r = (char *)cr8g;
|
||||
if (!r) {
|
||||
parseChordalRing(&cr8g);
|
||||
if (nGpu == 4 && is_xgmi) r = (char *)ring_4p3l;
|
||||
if (nGpu == 8 && cr8g) r = (char *)cr8g;
|
||||
if (nGpu == 8 && !cr8g) {
|
||||
r = (char *)ring_8p1h;
|
||||
if(!workgroups) workgroups = 16;
|
||||
}
|
||||
if (nGpu == 16) {
|
||||
r = (char *)ring_16p1h;
|
||||
if(!workgroups) workgroups = 24;
|
||||
}
|
||||
}
|
||||
|
||||
if(!workgroups) workgroups = 1;
|
||||
// clockwise and counter clockwise rings
|
||||
int ring[MAX_WORKGROUPS][MAX_GPU];
|
||||
for (int i = 0; i < MAX_WORKGROUPS; i++)
|
||||
@@ -433,14 +458,27 @@ int main(int argc,char* argv[])
|
||||
int num_rings = 0;
|
||||
if (r) {
|
||||
int j = 0, n = 0;
|
||||
int state = 0;
|
||||
do {
|
||||
if (r[n] == ' ') continue;
|
||||
if (r[n] == '|') {
|
||||
num_rings ++;
|
||||
j = 0;
|
||||
continue;
|
||||
int digit = r[n] - '0';
|
||||
if (digit >= 0 && digit <= 9) {
|
||||
if (state)
|
||||
ring[num_rings][j] = ring[num_rings][j]*10 + digit;
|
||||
else {
|
||||
ring[num_rings][j] = digit;
|
||||
state = 1;
|
||||
}
|
||||
}
|
||||
else {
|
||||
state = 0;
|
||||
j++;
|
||||
if (r[n] == ' ') continue;
|
||||
if (r[n] == '|') {
|
||||
num_rings ++;
|
||||
j = 0;
|
||||
continue;
|
||||
}
|
||||
}
|
||||
ring[num_rings][j++] = r[n] - '0';
|
||||
} while (r[n++] != 0x0);
|
||||
num_rings ++;
|
||||
} else {
|
||||
@@ -461,6 +499,10 @@ int main(int argc,char* argv[])
|
||||
struct profiling_data_t *profiling_data[MAX_GPU], *d_profiling_data[MAX_GPU];
|
||||
hipStream_t stream[MAX_GPU];
|
||||
|
||||
uint64_t *remOpCount, *d_remOpCount;
|
||||
HIPCHECK(hipHostMalloc((void**)&remOpCount, sizeof(uint64_t)*MAX_GPU, hipHostMallocMapped));
|
||||
HIPCHECK(hipHostGetDevicePointer((void**)&d_remOpCount, (void*)remOpCount, 0));
|
||||
|
||||
// print rings
|
||||
for (int i = 0; i < workgroups; i++) {
|
||||
printRing(i, ring[i], nGpu);
|
||||
@@ -519,16 +561,7 @@ int main(int argc,char* argv[])
|
||||
h_transfer_data[i].N = N;
|
||||
h_transfer_data[i].gpu = i;
|
||||
h_transfer_data[i].ngpu = nGpu;
|
||||
}
|
||||
|
||||
for (int i = 0; i < nGpu; i ++) {
|
||||
for (int j = 0; j < workgroups; j++) {
|
||||
for (int k = 0; k < nGpu; k++) {
|
||||
h_transfer_data[i].srcs[j][k] = buff[((i+k)%nGpu)*MAX_WORKGROUPS+j];
|
||||
h_transfer_data[i].dsts[j][k] = buff[((i+k)%nGpu)*MAX_WORKGROUPS+j] + N;
|
||||
//printf("Setup GPU %d bid %d srcs[%d] %p dsts[%d] %p\n", i, j, k, h_transfer_data[i].srcs[j][k], k, h_transfer_data[i].dsts[j][k]);
|
||||
}
|
||||
}
|
||||
h_transfer_data[i].remOpCount = d_remOpCount;
|
||||
}
|
||||
|
||||
for (int i = 0; i < nGpu; i ++) {
|
||||
@@ -543,13 +576,10 @@ int main(int argc,char* argv[])
|
||||
hipLaunchParams *launchParamsList= reinterpret_cast<hipLaunchParams *>(
|
||||
malloc(sizeof(hipLaunchParams)*MAX_GPU));
|
||||
|
||||
uint64_t opCount = 0;
|
||||
uint64_t opCount = workgroups;
|
||||
for (int op = begin_op; op < end_op; op ++) {
|
||||
if (op == OP_ALL2ALL && nGpu != 4 && nGpu != 8) {
|
||||
printf("\n%s only supports 4 or 8 GPUs.\n", ops[op]);
|
||||
continue;
|
||||
}
|
||||
printf("\n[Testing %s]: \n", ops[op]);
|
||||
const char *OpsName[] = {"Copy", "Local Copy", "Double Copy", "doublecopylocal", "Reduce", "ReduceCopy", "Read"};
|
||||
printf("\n[Testing %s]: \n", OpsName[op]);
|
||||
// 4 warm up cycles
|
||||
for (int j = 0; j < 4; j ++) {
|
||||
for (int i = 0; i < nGpu; i ++) {
|
||||
@@ -557,10 +587,8 @@ int main(int argc,char* argv[])
|
||||
args[i*3] = &transfer_data[i];
|
||||
args[i*3+1] = &d_profiling_data[i];
|
||||
args[i*3+2] = &opCount;
|
||||
if (op == OP_ALL2ALL)
|
||||
launchParamsList[i].func = reinterpret_cast<void *>(flagSyncKerns[op + (nGpu/8)]);
|
||||
else
|
||||
launchParamsList[i].func = reinterpret_cast<void *>(flagSyncKerns[op]);
|
||||
launchParamsList[i].func =
|
||||
reinterpret_cast<void *>(flagSyncKerns[op*2 + sync]);
|
||||
launchParamsList[i].gridDim = dim3(workgroups, 1, 1),
|
||||
launchParamsList[i].blockDim = dim3(THREADS, 1, 1),
|
||||
launchParamsList[i].sharedMem = 0;
|
||||
@@ -572,7 +600,7 @@ int main(int argc,char* argv[])
|
||||
#else
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
//launch the kernel
|
||||
hipLaunchKernelGGL(flagSyncKerns[op == OP_ALL2ALL ? op + (nGpu/8) : op],
|
||||
hipLaunchKernelGGL(flagSyncKerns[op*2 + sync],
|
||||
/*grid dim x,y,z*/ dim3(workgroups, 1, 1),
|
||||
/*block dim x,y,z*/ dim3(THREADS, 1, 1),
|
||||
/*dynamic shared mem*/ 0,
|
||||
@@ -580,7 +608,7 @@ int main(int argc,char* argv[])
|
||||
/*kernel args*/ transfer_data[i], d_profiling_data[i], opCount);
|
||||
}
|
||||
#endif
|
||||
opCount++;
|
||||
opCount+=workgroups;
|
||||
}
|
||||
|
||||
for (int i = 0; i < nGpu; i ++) {
|
||||
@@ -596,10 +624,8 @@ int main(int argc,char* argv[])
|
||||
args[i*3] = &transfer_data[i];
|
||||
args[i*3+1] = &d_profiling_data[i];
|
||||
args[i*3+2] = &opCount;
|
||||
if (op == OP_ALL2ALL)
|
||||
launchParamsList[i].func = reinterpret_cast<void *>(flagSyncKerns[op + (nGpu/8)]);
|
||||
else
|
||||
launchParamsList[i].func = reinterpret_cast<void *>(flagSyncKerns[op]);
|
||||
launchParamsList[i].func =
|
||||
reinterpret_cast<void *>(flagSyncKerns[op*2 + sync]);
|
||||
launchParamsList[i].gridDim = dim3(workgroups, 1, 1),
|
||||
launchParamsList[i].blockDim = dim3(THREADS, 1, 1),
|
||||
launchParamsList[i].sharedMem = 0;
|
||||
@@ -611,7 +637,7 @@ int main(int argc,char* argv[])
|
||||
#else
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
//launch the kernel
|
||||
hipLaunchKernelGGL(flagSyncKerns[op == OP_ALL2ALL ? op + (nGpu/8) : op],
|
||||
hipLaunchKernelGGL(flagSyncKerns[op*2 + sync],
|
||||
/*grid dim x,y,z*/ dim3(workgroups, 1, 1),
|
||||
/*block dim x,y,z*/ dim3(THREADS, 1, 1),
|
||||
/*dynamic shared mem*/ 0,
|
||||
@@ -619,7 +645,7 @@ int main(int argc,char* argv[])
|
||||
/*kernel args*/ transfer_data[i], d_profiling_data[i], opCount);
|
||||
}
|
||||
#endif
|
||||
opCount++;
|
||||
opCount+=workgroups;
|
||||
}
|
||||
|
||||
for (int i = 0; i < nGpu; i ++) {
|
||||
@@ -650,46 +676,24 @@ int main(int argc,char* argv[])
|
||||
uint32_t hopcount;
|
||||
HIPCHECK(hipExtGetLinkTypeAndHopCount(i, next_gpu , &linktype, &hopcount));
|
||||
|
||||
if (op == OP_ALL2ALL) {
|
||||
if(prop.gcnArch == 906) {
|
||||
write_cycle = write_cycle + profiling_data[i]->write_cycles[j];
|
||||
bytes_transferred = bytes_transferred + profiling_data[i]->bytes_transferred[j];
|
||||
double t0 = (double)profiling_data[i]->write_cycles[j]/RTC_CLOCK_FREQ_VEGA20;
|
||||
fprintf(stderr, "%-20d %-d<->all %-13d %-13s %-13.4f %-20lu %-.2f\n",
|
||||
i, i, j, link_type_name[linktype], t0, profiling_data[i]->bytes_transferred[j], (double)profiling_data[i]->bytes_transferred[j]/(t0*1.0E9));
|
||||
} else if (prop.gcnArch == 908) {
|
||||
write_cycle = write_cycle + profiling_data[i]->write_cycles[j];
|
||||
bytes_transferred = bytes_transferred + profiling_data[i]->bytes_transferred[j];
|
||||
double t0 = (double)profiling_data[i]->write_cycles[j]/RTC_CLOCK_FREQ_ARCTURUS;
|
||||
fprintf(stderr, "%-20d %-d<->all %-13d %-13s %-13.4f %-20lu %-.2f\n",
|
||||
i, i, j, link_type_name[linktype], t0, profiling_data[i]->bytes_transferred[j], (double)profiling_data[i]->bytes_transferred[j]/(t0*1.0E9));
|
||||
} else {
|
||||
write_cycle = write_cycle + profiling_data[i]->write_cycles[j];
|
||||
bytes_transferred = bytes_transferred + profiling_data[i]->bytes_transferred[j];
|
||||
double t0 = (double)profiling_data[i]->write_cycles[j]/RTC_CLOCK_FREQ_DEFAULT;
|
||||
fprintf(stderr, "%-20d %-d<->all %-13d %-13s %-13.4f %-20lu %-.2f\n",
|
||||
i, i, j, link_type_name[linktype], t0, profiling_data[i]->bytes_transferred[j], (double)profiling_data[i]->bytes_transferred[j]/(t0*1.0E9));
|
||||
}
|
||||
if(prop.gcnArch == 906) {
|
||||
write_cycle = write_cycle + profiling_data[i]->write_cycles[j];
|
||||
bytes_transferred = bytes_transferred + profiling_data[i]->bytes_transferred[j];
|
||||
double t0 = (double)profiling_data[i]->write_cycles[j]/((double)RTC_CLOCK_FREQ_VEGA20);
|
||||
fprintf(stderr, "%-20d %-d->%-10d %-13d %-13s %-13.4f %-20lu %-.2f\n",
|
||||
i,i, next_gpu,j,link_type_name[linktype],t0, profiling_data[i]->bytes_transferred[j], (double)profiling_data[i]->bytes_transferred[j]/(t0*1.0E9));
|
||||
} else if (prop.gcnArch == 908) {
|
||||
write_cycle = write_cycle + profiling_data[i]->write_cycles[j];
|
||||
bytes_transferred = bytes_transferred + profiling_data[i]->bytes_transferred[j];
|
||||
double t0 = (double)profiling_data[i]->write_cycles[j]/((double)RTC_CLOCK_FREQ_ARCTURUS);
|
||||
fprintf(stderr, "%-20d %-d->%-10d %-13d %-13s %-13.4f %-20lu %-.2f\n",
|
||||
i,i, next_gpu,j,link_type_name[linktype],t0, profiling_data[i]->bytes_transferred[j], (double)profiling_data[i]->bytes_transferred[j]/(t0*1.0E9));
|
||||
} else {
|
||||
if(prop.gcnArch == 906) {
|
||||
write_cycle = write_cycle + profiling_data[i]->write_cycles[j];
|
||||
bytes_transferred = bytes_transferred + profiling_data[i]->bytes_transferred[j];
|
||||
double t0 = (double)profiling_data[i]->write_cycles[j]/RTC_CLOCK_FREQ_VEGA20;
|
||||
fprintf(stderr, "%-20d %-d->%-10d %-13d %-13s %-13.4f %-20lu %-.2f\n",
|
||||
i, i, next_gpu, j, link_type_name[linktype], t0, profiling_data[i]->bytes_transferred[j], (double)profiling_data[i]->bytes_transferred[j]/(t0*1.0E9));
|
||||
} else if (prop.gcnArch == 908) {
|
||||
write_cycle = write_cycle + profiling_data[i]->write_cycles[j];
|
||||
bytes_transferred = bytes_transferred + profiling_data[i]->bytes_transferred[j];
|
||||
double t0 = (double)profiling_data[i]->write_cycles[j]/RTC_CLOCK_FREQ_ARCTURUS;
|
||||
fprintf(stderr, "%-20d %-d->%-10d %-13d %-13s %-13.4f %-20lu %-.2f\n",
|
||||
i, i, next_gpu, j, link_type_name[linktype], t0, profiling_data[i]->bytes_transferred[j], (double)profiling_data[i]->bytes_transferred[j]/(t0*1.0E9));
|
||||
} else {
|
||||
write_cycle = write_cycle + profiling_data[i]->write_cycles[j];
|
||||
bytes_transferred = bytes_transferred + profiling_data[i]->bytes_transferred[j];
|
||||
double t0 = (double)profiling_data[i]->write_cycles[j]/RTC_CLOCK_FREQ_DEFAULT;
|
||||
fprintf(stderr, "%-20d %-d->%-10d %-13d %-13s %-13.4f %-20lu %-.2f\n",
|
||||
i, i, next_gpu, j, link_type_name[linktype], t0, profiling_data[i]->bytes_transferred[j], (double)profiling_data[i]->bytes_transferred[j]/(t0*1.0E9));
|
||||
}
|
||||
write_cycle = write_cycle + profiling_data[i]->write_cycles[j];
|
||||
bytes_transferred = bytes_transferred + profiling_data[i]->bytes_transferred[j];
|
||||
double t0 = (double)profiling_data[i]->write_cycles[j]/((double)RTC_CLOCK_FREQ_DEFAULT);
|
||||
fprintf(stderr, "%-20d %-d->%-10d %-13d %-13s %-13.4f %-20lu %-.2f\n",
|
||||
i,i, next_gpu,j,link_type_name[linktype],t0, profiling_data[i]->bytes_transferred[j], (double)profiling_data[i]->bytes_transferred[j]/(t0*1.0E9));
|
||||
}
|
||||
}
|
||||
print_table_summary_line();
|
||||
@@ -724,4 +728,10 @@ int main(int argc,char* argv[])
|
||||
HIPCHECK(hipFree((void*) d_profiling_data[i]));
|
||||
free(profiling_data[i]);
|
||||
}
|
||||
|
||||
printf("opCount: ");
|
||||
for (int i = 0; i < nGpu; i++)
|
||||
printf("%ld ", remOpCount[i]);
|
||||
printf("\n");
|
||||
HIPCHECK(hipHostFree((void*)remOpCount));
|
||||
}
|
||||
|
||||
Reference in New Issue
Block a user