diff --git a/projects/rccl/tools/rccl-prim-test/rccl_prim_test.cpp b/projects/rccl/tools/rccl-prim-test/rccl_prim_test.cpp index 5462c1031c..b63fd058eb 100644 --- a/projects/rccl/tools/rccl-prim-test/rccl_prim_test.cpp +++ b/projects/rccl/tools/rccl-prim-test/rccl_prim_test.cpp @@ -33,6 +33,7 @@ THE SOFTWARE. #include #include "copy_kernel.h" +#define MAX_GPU 8 #define MAX_WORKGROUPS 8 #define THREADS 256 @@ -42,12 +43,14 @@ THE SOFTWARE. #define REDUCECOPY_UNROLL 2 struct transfer_data_t { - float *dest0; //remote fine grain - float *src0; //local fine grain - float *dest1; //local coarse grain - float *src1; //local coarse grain + float *dest0[MAX_WORKGROUPS]; //remote fine grain + float *src0[MAX_WORKGROUPS]; //local fine grain + float *dest1[MAX_WORKGROUPS]; //local coarse grain + float *src1[MAX_WORKGROUPS]; //local coarse grain int N; int gpu; + int ngpu; + uint64_t *remOpCount; }; struct profiling_data_t { @@ -68,22 +71,34 @@ enum Ops { NUM_OPS, }; -template -__global__ void flag_sync_kernel(struct transfer_data_t* transfer_data, struct profiling_data_t* profiling_data) { +template +__global__ void flag_sync_kernel(struct transfer_data_t* transfer_data, struct profiling_data_t* profiling_data, uint64_t opCount) { size_t idx = threadIdx.x; uint64_t curr_time, next_time; + int bid = blockIdx.x; + int n = transfer_data->N; + + // signal self ready and wait until all GPUs are ready + if (idx == 0) { + if (bid == 0) + STORE(&transfer_data->remOpCount[transfer_data->gpu], opCount); + if (sync) { + for (int i = 0; i < transfer_data->ngpu; i++) { + while (LOAD(&transfer_data->remOpCount[i]) < opCount) {}; + } + } + } + __syncthreads(); if (idx == 0) { curr_time = clock64(); } - int offset = transfer_data->N * blockIdx.x / gridDim.x; - int n = transfer_data->N / gridDim.x; - if (op == OP_COPY) Copy(transfer_data->dest0 + offset, transfer_data->src0 + offset, n); - if (op == OP_LOCALCOPY) Copy(transfer_data->dest1 + offset, transfer_data->src0 + offset, n); - if (op == OP_DOUBLECOPY) DoubleCopy(transfer_data->dest0 + offset, transfer_data->dest1 + offset, transfer_data->src0 + offset, n); - if (op == OP_REDUCE) Reduce(transfer_data->dest0 + offset, transfer_data->src0 + offset, transfer_data->src1 + offset, n); - if (op == OP_REDUCECOPY) ReduceCopy(transfer_data->dest0 + offset, transfer_data->dest1 + offset, transfer_data->src0 + offset, transfer_data->src1 + offset, n); + if (op == OP_COPY) Copy(transfer_data->dest0[bid], transfer_data->src0[bid], n); + if (op == OP_LOCALCOPY) Copy(transfer_data->dest1[bid], transfer_data->src0[bid], n); + if (op == OP_DOUBLECOPY) DoubleCopy(transfer_data->dest0[bid], transfer_data->dest1[bid], transfer_data->src0[bid], n); + if (op == OP_REDUCE) Reduce(transfer_data->dest0[bid], transfer_data->src0[bid], transfer_data->src1[bid], n); + if (op == OP_REDUCECOPY) ReduceCopy(transfer_data->dest0[bid], transfer_data->dest1[bid], transfer_data->src0[bid], transfer_data->src1[bid], n); __syncthreads(); if (idx == 0) { @@ -93,14 +108,19 @@ __global__ void flag_sync_kernel(struct transfer_data_t* transfer_data, struct p } } -typedef void(*flag_sync_kernel_t)(struct transfer_data_t* transfer_data, struct profiling_data_t* profiling_data); +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] = { - flag_sync_kernel, - flag_sync_kernel, - flag_sync_kernel, - flag_sync_kernel, - flag_sync_kernel, +static flag_sync_kernel_t const flagSyncKerns[NUM_OPS*2] = { + flag_sync_kernel, + flag_sync_kernel, + flag_sync_kernel, + flag_sync_kernel, + flag_sync_kernel, + flag_sync_kernel, + flag_sync_kernel, + flag_sync_kernel, + flag_sync_kernel, + flag_sync_kernel, }; __global__ void initTestDataKernel(float* data, const size_t N, const int gpu) { @@ -122,21 +142,83 @@ do { \ } \ } while (0) -static void setupPeers() { - int deviceCnt, dev; +static void setupPeers(uint32_t *info) { + int deviceCnt, dev; - HIPCHECK(hipGetDeviceCount(&deviceCnt)); - HIPCHECK(hipGetDevice(&dev)); - //! If gpus are not peer enabled, enable them - for (int i = 0; i < deviceCnt; i++) { - HIPCHECK(hipSetDevice(i)); - for (int j = 0; j < deviceCnt; j++) { - if (i != j) { - HIPCHECK(hipDeviceEnablePeerAccess(j, 0)); - } + HIPCHECK(hipGetDeviceCount(&deviceCnt)); + HIPCHECK(hipGetDevice(&dev)); + //! If gpus are not peer enabled, enable them + for (int i = 0; i < deviceCnt; i++) { + HIPCHECK(hipSetDevice(i)); + for (int j = 0; j < deviceCnt; j++) { + if (i != j) { + 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); + exit(-1); } + HIPCHECK(hipDeviceEnablePeerAccess(j, 0)); + uint32_t linktype; + HIPCHECK(hipExtGetLinkTypeAndHopCount(i, j, &linktype, &info[i*deviceCnt+j])); + } + else + info[i*deviceCnt+j] = 0; } - HIPCHECK(hipSetDevice(dev)); + } + HIPCHECK(hipSetDevice(dev)); +} + +static void printRing(int id, int *ring, int deviceCnt) { + printf("Ring %d: ", id); + for (int i = 0; i < deviceCnt; i++) + printf("%1d ", ring[i]); + printf("\n"); +} + +static void findConnect(uint32_t *info, int *ring, int deviceCnt) { + int n = 0, curr = 0, best; + uint32_t temp[MAX_GPU*MAX_GPU]; + for (int i = 0; i < deviceCnt*deviceCnt; i++) temp[i] = 0; + for (int i = 0; i < deviceCnt; i++) { + for (int j = 0; j < deviceCnt; j++) temp[j*deviceCnt+curr] = 1; + ring[n] = curr; + n++; + int hops = 99; + for (int j = 0; j < deviceCnt; j++) { + if (temp[curr*deviceCnt+j]) continue; + if (info[curr*deviceCnt+j] < hops) { + best = j; + hops = info[curr*deviceCnt+j]; + } + } + curr = best; + } +} + +static int findNextGpu(int *ring, int gpu, int deviceCnt) { + int i; + for (i = 0; i < deviceCnt; i ++) + if (ring[i] == gpu) break; + return ring[(i+1)%deviceCnt]; +} + +static void setupRings(uint32_t *info, int *ring_0, int *ring_1) { + int deviceCnt, dev; + HIPCHECK(hipGetDeviceCount(&deviceCnt)); + printf("Connection matrix:\n"); + for (int i = 0; i < deviceCnt; i++) { + for (int j = 0; j < deviceCnt; j++) + printf("%2d ", info[i*deviceCnt+j]); + printf("\n"); + } + findConnect(info, ring_0, deviceCnt); + printRing(0, ring_0, deviceCnt); + ring_1[0] =0; + for (int i = 1; i < deviceCnt; i++) + ring_1[i] = ring_0[deviceCnt-i]; + printRing(1, ring_1, deviceCnt); } char* getCmdOption(char ** begin, char ** end, const std::string & option) { @@ -155,7 +237,7 @@ bool cmdOptionExists(char** begin, char** end, const std::string& option) { int main(int argc,char* argv[]) { if (cmdOptionExists(argv, argv + argc, "-h")) { - printf("./rccl_prim_test -w num_workgroups -p copy|localcopy|doublecopy|reduce|reducecopy|all -n iterations\n"); + printf("./rccl_prim_test -w num_workgroups -p copy|localcopy|doublecopy|reduce|reducecopy|all -i iterations -n bytes -s 0|1\n"); exit(0); } @@ -166,11 +248,24 @@ int main(int argc,char* argv[]) printf("Benchmarking using %d workgroups\n", workgroups); int iters = 10; - char *it = getCmdOption(argv, argv + argc, "-n"); + char *it = getCmdOption(argv, argv + argc, "-i"); if (it) iters = atol(it); printf("Benchmarking using %d iterations\n", iters); + uint64_t nBytes = 2097152; + char *nb = getCmdOption(argv, argv + argc, "-n"); + if (nb) + nBytes = atol(nb); + printf("Benchmarking using %ld bytes\n", nBytes); + uint64_t N = nBytes/sizeof(float); + + int sync = 0; + char *s = getCmdOption(argv, argv + argc, "-s"); + if (s) + sync = atol(s); + if (sync) printf("Sync all GPUs before operation\n"); + const char *ops[] = {"copy", "localcopy", "doublecopy", "reduce", "reducecopy", "all"}; char *prim = getCmdOption(argv, argv + argc, "-p"); int op = 5, begin_op, end_op; @@ -188,194 +283,162 @@ int main(int argc,char* argv[]) printf("Benchmarking all ops\n"); } + uint32_t connection_info[MAX_GPU*MAX_GPU]; // Enable peer access - setupPeers(); + setupPeers(connection_info); + // clockwise and counter clockwise rings + int ring_0[MAX_GPU] = {-1, -1, -1, -1}; + int ring_1[MAX_GPU] = {-1, -1, -1, -1}; + setupRings(connection_info, ring_0, ring_1); // data buffers - float *buff_0, *buff_1, *buff_coarse_0, *buff_coarse_1; - struct transfer_data_t h_transfer_data_0, h_transfer_data_1, *transfer_data_0, *transfer_data_1; - struct profiling_data_t *profiling_data_0, *profiling_data_1, *d_profiling_data_0, *d_profiling_data_1; - uint64_t N = 2097152*4*MAX_WORKGROUPS; + float *buff[MAX_GPU*MAX_WORKGROUPS], *buff_coarse[MAX_GPU*MAX_WORKGROUPS]; + struct transfer_data_t h_transfer_data[MAX_GPU], *transfer_data[MAX_GPU]; + struct profiling_data_t *profiling_data[MAX_GPU], *d_profiling_data[MAX_GPU]; + hipStream_t stream[MAX_GPU]; - int hipDev = 0; - HIPCHECK(hipSetDevice(hipDev)); - hipDeviceProp_t prop; - HIPCHECK(hipGetDeviceProperties(&prop, hipDev)); - printf("# device %d [0x%02x] %s\n", - hipDev, prop.pciBusID, prop.name); - HIPCHECK(hipExtMallocWithFlags((void**) &transfer_data_0, sizeof(struct transfer_data_t), hipDeviceMallocFinegrained)); - //printf("GPU 0: allocated fine grain VRAM at %llx\n", (unsigned long long)transfer_data_0); - HIPCHECK(hipExtMallocWithFlags((void**) &buff_0, 2*N*sizeof(float), hipDeviceMallocFinegrained)); - //printf("GPU 0: allocated fine grain VRAM at %llx\n", (unsigned long long)buff_0); - HIPCHECK(hipMalloc((void**) &buff_coarse_0, 2*N*sizeof(float))); - //printf("GPU 0: allocated coarse grain VRAM at %llx\n", (unsigned long long)buff_coarse_0); - profiling_data_0 = (struct profiling_data_t *)malloc(sizeof(struct profiling_data_t)); - HIPCHECK(hipMalloc((void**) &d_profiling_data_0, sizeof(struct profiling_data_t))); - //create stream - hipStream_t stream_0; - HIPCHECK(hipStreamCreate(&stream_0)); - //randomize test data - hipLaunchKernelGGL(initTestDataKernel, - /*grid dim x,y,z*/ dim3(32, 1, 1), - /*block dim x,y,z*/ dim3(THREADS, 1, 1), - /*dynamic shared mem*/ 0, - /*stream*/ stream_0, - /*kernel args*/ buff_0, 2*N, 0); - hipLaunchKernelGGL(initTestDataKernel, - /*grid dim x,y,z*/ dim3(32, 1, 1), - /*block dim x,y,z*/ dim3(THREADS, 1, 1), - /*dynamic shared mem*/ 0, - /*stream*/ stream_0, - /*kernel args*/ buff_coarse_0, 2*N, 0); + int nGpu = 1; + HIPCHECK(hipGetDeviceCount(&nGpu)); + uint64_t *remOpCount, *d_remOpCount; + HIPCHECK(hipHostMalloc((void**)&remOpCount, sizeof(uint64_t)*MAX_GPU, hipHostMallocMapped)); + HIPCHECK(hipHostGetDevicePointer((void**)&d_remOpCount, (void*)remOpCount, 0)); - hipDev = 1; - HIPCHECK(hipSetDevice(hipDev)); - HIPCHECK(hipGetDeviceProperties(&prop, hipDev)); - printf("# device %d [0x%02x] %s\n", - hipDev, prop.pciBusID, prop.name); - HIPCHECK(hipExtMallocWithFlags((void**) &transfer_data_1, sizeof(struct transfer_data_t), hipDeviceMallocFinegrained)); - //printf("GPU 1: allocated fine grain VRAM at %llx\n", (unsigned long long)transfer_data_1); - HIPCHECK(hipExtMallocWithFlags((void**) &buff_1, 2*N*sizeof(float), hipDeviceMallocFinegrained)); - //printf("GPU 1: allocated fine grain VRAM at %llx\n", (unsigned long long)buff_1); - HIPCHECK(hipMalloc((void**) &buff_coarse_1, 2*N*sizeof(float))); - //printf("GPU 1: allocated coarse grain VRAM at %llx\n", (unsigned long long)buff_coarse_1); - profiling_data_1 = (struct profiling_data_t *)malloc(sizeof(struct profiling_data_t)); - HIPCHECK(hipMalloc((void**) &d_profiling_data_1, sizeof(struct profiling_data_t))); - //create stream - hipStream_t stream_1; - HIPCHECK(hipStreamCreate(&stream_1)); - //randomize test data - hipLaunchKernelGGL(initTestDataKernel, - /*grid dim x,y,z*/ dim3(32, 1, 1), - /*block dim x,y,z*/ dim3(THREADS, 1, 1), - /*dynamic shared mem*/ 0, - /*stream*/ stream_1, - /*kernel args*/ buff_1, 2*N, 1); - hipLaunchKernelGGL(initTestDataKernel, - /*grid dim x,y,z*/ dim3(32, 1, 1), - /*block dim x,y,z*/ dim3(THREADS, 1, 1), - /*dynamic shared mem*/ 0, - /*stream*/ stream_1, - /*kernel args*/ buff_coarse_1, 2*N, 1); - h_transfer_data_0.dest0 = buff_1; - h_transfer_data_0.dest1 = buff_coarse_0 + N; - h_transfer_data_0.src0 = buff_0; - h_transfer_data_0.src1 = buff_coarse_0; - h_transfer_data_0.N = N; - h_transfer_data_0.gpu = 0; + for (int i = 0; i < nGpu; i ++) { + HIPCHECK(hipSetDevice(i)); + hipDeviceProp_t prop; + HIPCHECK(hipGetDeviceProperties(&prop, i)); + printf("# device %d [0x%02x] %s\n", + i, prop.pciBusID, prop.name); + //create stream + HIPCHECK(hipStreamCreate(&stream[i])); + profiling_data[i] = (struct profiling_data_t *)malloc(sizeof(struct profiling_data_t)); + HIPCHECK(hipMalloc((void**) &d_profiling_data[i], sizeof(struct profiling_data_t))); - h_transfer_data_1.dest0 = buff_0 + N; - h_transfer_data_1.dest1 = buff_coarse_1; - h_transfer_data_1.src0 = buff_1 + N; - h_transfer_data_1.src1 = buff_coarse_1 + N; - h_transfer_data_1.N = N; - h_transfer_data_1.gpu = 1; + HIPCHECK(hipExtMallocWithFlags((void**) &transfer_data[i], sizeof(struct transfer_data_t), hipDeviceMallocFinegrained)); + for (int j = 0; j < workgroups; j++) { + HIPCHECK(hipExtMallocWithFlags((void**) &buff[i*MAX_WORKGROUPS+j], 2*N*sizeof(float), hipDeviceMallocFinegrained)); + HIPCHECK(hipMalloc((void**) &buff_coarse[i*MAX_WORKGROUPS+j], 2*N*sizeof(float))); + //randomize test data + hipLaunchKernelGGL(initTestDataKernel, + /*grid dim x,y,z*/ dim3(32, 1, 1), + /*block dim x,y,z*/ dim3(THREADS, 1, 1), + /*dynamic shared mem*/ 0, + /*stream*/ stream[i], + /*kernel args*/ buff[i*MAX_WORKGROUPS+j], 2*N, 0); + hipLaunchKernelGGL(initTestDataKernel, + /*grid dim x,y,z*/ dim3(32, 1, 1), + /*block dim x,y,z*/ dim3(THREADS, 1, 1), + /*dynamic shared mem*/ 0, + /*stream*/ stream[i], + /*kernel args*/ buff_coarse[i*MAX_WORKGROUPS+j], 2*N, 0); + } + } - HIPCHECK(hipSetDevice(0)); - HIPCHECK(hipMemcpyAsync(transfer_data_0, &h_transfer_data_0, - sizeof(struct transfer_data_t), hipMemcpyHostToDevice, - stream_0)); - HIPCHECK(hipStreamSynchronize(stream_0)); + for (int i = 0; i < nGpu; i ++) { + for (int j = 0; j < workgroups; j++) { + int next_gpu; + if (j%2) + next_gpu = findNextGpu(ring_1, i, nGpu); + else + next_gpu = findNextGpu(ring_0, i, nGpu); + //printf("GPU %d Ring %d -> Next GPU %d\n", i, j, next_gpu); + h_transfer_data[i].dest0[j] = buff[next_gpu*MAX_WORKGROUPS+j] + N; + h_transfer_data[i].dest1[j] = buff_coarse[i*MAX_WORKGROUPS+j] + N; + h_transfer_data[i].src0[j] = buff[i*MAX_WORKGROUPS+j]; + h_transfer_data[i].src1[j] = buff_coarse[i*MAX_WORKGROUPS+j]; + } + h_transfer_data[i].N = N; + h_transfer_data[i].gpu = i; + h_transfer_data[i].ngpu = nGpu; + h_transfer_data[i].remOpCount = d_remOpCount; + } - HIPCHECK(hipSetDevice(1)); - HIPCHECK(hipMemcpyAsync(transfer_data_1, &h_transfer_data_1, - sizeof(struct transfer_data_t), hipMemcpyHostToDevice, - stream_1)); - HIPCHECK(hipStreamSynchronize(stream_1)); + for (int i = 0; i < nGpu; i ++) { + HIPCHECK(hipSetDevice(i)); + HIPCHECK(hipMemcpyAsync(transfer_data[i], &h_transfer_data[i], + sizeof(struct transfer_data_t), hipMemcpyHostToDevice, + stream[i])); + HIPCHECK(hipStreamSynchronize(stream[i])); + } + uint64_t opCount = 0; for (int op = begin_op; op < end_op; op ++) { const char *OpsName[] = {"Copy", "Local Copy", "Double Copy", "Reduce", "ReduceCopy"}; printf("Testing %s: \n", OpsName[op]); // 2 warm up cycles for (int i = 0; i < 2; i ++) { - HIPCHECK(hipSetDevice(0)); - //launch the kernel - hipLaunchKernelGGL(flagSyncKerns[op], - /*grid dim x,y,z*/ dim3(workgroups, 1, 1), - /*block dim x,y,z*/ dim3(THREADS, 1, 1), - /*dynamic shared mem*/ 0, - /*stream*/ stream_0, - /*kernel args*/ transfer_data_0, d_profiling_data_0); - - HIPCHECK(hipSetDevice(1)); - //launch the kernel - hipLaunchKernelGGL(flagSyncKerns[op], - /*grid dim x,y,z*/ dim3(workgroups, 1, 1), - /*block dim x,y,z*/ dim3(THREADS, 1, 1), - /*dynamic shared mem*/ 0, - /*stream*/ stream_1, - /*kernel args*/ transfer_data_1, d_profiling_data_1); + for (int i = 0; i < nGpu; i ++) { + HIPCHECK(hipSetDevice(i)); + //launch the kernel + 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, + /*stream*/ stream[i], + /*kernel args*/ transfer_data[i], d_profiling_data[i], opCount); + } + opCount++; } - HIPCHECK(hipSetDevice(0)); - HIPCHECK(hipStreamSynchronize(stream_0)); - HIPCHECK(hipMemset(d_profiling_data_0, 0, sizeof(struct profiling_data_t))); - HIPCHECK(hipSetDevice(1)); - HIPCHECK(hipStreamSynchronize(stream_1)); - HIPCHECK(hipMemset(d_profiling_data_1, 0, sizeof(struct profiling_data_t))); + for (int i = 0; i < nGpu; i ++) { + HIPCHECK(hipSetDevice(i)); + HIPCHECK(hipStreamSynchronize(stream[i])); + HIPCHECK(hipMemset(d_profiling_data[i], 0, sizeof(struct profiling_data_t))); + } auto start = std::chrono::high_resolution_clock::now(); for (int i = 0; i < iters; i ++) { - HIPCHECK(hipSetDevice(0)); - //launch the kernel - hipLaunchKernelGGL(flagSyncKerns[op], - /*grid dim x,y,z*/ dim3(workgroups, 1, 1), - /*block dim x,y,z*/ dim3(THREADS, 1, 1), - /*dynamic shared mem*/ 0, - /*stream*/ stream_0, - /*kernel args*/ transfer_data_0, d_profiling_data_0); - - HIPCHECK(hipSetDevice(1)); - //launch the kernel - hipLaunchKernelGGL(flagSyncKerns[op], - /*grid dim x,y,z*/ dim3(workgroups, 1, 1), - /*block dim x,y,z*/ dim3(THREADS, 1, 1), - /*dynamic shared mem*/ 0, - /*stream*/ stream_1, - /*kernel args*/ transfer_data_1, d_profiling_data_1); + for (int i = 0; i < nGpu; i ++) { + HIPCHECK(hipSetDevice(i)); + //launch the kernel + 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, + /*stream*/ stream[i], + /*kernel args*/ transfer_data[i], d_profiling_data[i], opCount); + } + opCount++; + } + + for (int i = 0; i < nGpu; i ++) { + HIPCHECK(hipSetDevice(i)); + HIPCHECK(hipStreamSynchronize(stream[i])); } - HIPCHECK(hipSetDevice(0)); - HIPCHECK(hipStreamSynchronize(stream_0)); - HIPCHECK(hipSetDevice(1)); - HIPCHECK(hipStreamSynchronize(stream_1)); auto delta = std::chrono::high_resolution_clock::now() - start; double deltaSec = std::chrono::duration_cast>(delta).count(); - HIPCHECK(hipMemcpyAsync(profiling_data_0, d_profiling_data_0, - sizeof(struct profiling_data_t), hipMemcpyDeviceToHost, - stream_0)); - HIPCHECK(hipStreamSynchronize(stream_0)); - - HIPCHECK(hipMemcpyAsync(profiling_data_1, d_profiling_data_1, - sizeof(struct profiling_data_t), hipMemcpyDeviceToHost, - stream_1)); - HIPCHECK(hipStreamSynchronize(stream_1)); - - double speed = (double)(profiling_data_0->bytes_transferred) / (deltaSec*1.0E9); - printf("Transfered %lu bytes in %f s. Throughput %f GB/s\n", profiling_data_0->bytes_transferred, deltaSec, speed); - + for (int i = 0; i < nGpu; i ++) { + HIPCHECK(hipMemcpyAsync(profiling_data[i], d_profiling_data[i], + sizeof(struct profiling_data_t), hipMemcpyDeviceToHost, + stream[i])); + HIPCHECK(hipStreamSynchronize(stream[i])); #define RTC_CLOCK_FREQ 2.7E07 - double t0 = (double)profiling_data_0->write_cycles/((double)RTC_CLOCK_FREQ)/(double)workgroups; - fprintf(stderr, "GPU 0: time %.4fs bytes_transferred %lu kernel throughput %.2f GB/s\n", - t0, profiling_data_0->bytes_transferred, (double)profiling_data_0->bytes_transferred/(t0*1.0E9)); + double t0 = (double)profiling_data[i]->write_cycles/((double)RTC_CLOCK_FREQ)/(double)workgroups; + fprintf(stderr, "GPU %d: time %.4fs bytes_transferred %lu kernel throughput %.2f GB/s\n", + i, t0, profiling_data[i]->bytes_transferred, (double)profiling_data[i]->bytes_transferred/(t0*1.0E9)); + } - double t1 = (double)profiling_data_1->write_cycles/((double)RTC_CLOCK_FREQ)/(double)workgroups; - fprintf(stderr, "GPU 1: time %.4fs bytes_transferred %lu kernel throughput %.2f GB/s\n", - t1, profiling_data_1->bytes_transferred, (double)profiling_data_0->bytes_transferred/(t1*1.0E9)); + double speed = (double)(profiling_data[0]->bytes_transferred) / (deltaSec*1.0E9); + printf("Transfered %lu bytes in %f s. Throughput %f GB/s\n", profiling_data[0]->bytes_transferred, deltaSec, speed); } - HIPCHECK(hipStreamDestroy(stream_0)); - HIPCHECK(hipStreamDestroy(stream_1)); - HIPCHECK(hipFree((void*) transfer_data_0)); - HIPCHECK(hipFree((void*) buff_0)); - HIPCHECK(hipFree((void*) buff_coarse_0)); - HIPCHECK(hipFree((void*) d_profiling_data_0)); - free(profiling_data_0); - HIPCHECK(hipFree((void*) transfer_data_1)); - HIPCHECK(hipFree((void*) buff_1)); - HIPCHECK(hipFree((void*) buff_coarse_1)); - HIPCHECK(hipFree((void*) d_profiling_data_1)); - free(profiling_data_1); + for (int i = 0; i < nGpu; i ++) { + HIPCHECK(hipStreamDestroy(stream[i])); + HIPCHECK(hipFree((void*) transfer_data[i])); + for (int j = 0; j < workgroups; j++) { + HIPCHECK(hipFree((void*) buff[i*MAX_WORKGROUPS+j])); + HIPCHECK(hipFree((void*) buff_coarse[i*MAX_WORKGROUPS+j])); + } + 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)); }