AI/rocm-systems
2
0
Garpu 0
Kode Masalah Tarik Permintaan Actions Packages Projects Rilis Wiki Kegiatan

Refactor primitive test to support multiple GPUs in rings (#94)

Telusuri Sumber 
* Refactor primitive test to support multiple GPUs in rings

* Make GPUs sync before transfer optional

* Use same ring format as RCCL

* Extend to 8 GPUs and report errors if there is no P2P access

* Control GPUs sync before ops from command line with "-s" option

* Change buffer size through command line option "-n"

Rename iterations command line option to "-i"


[ROCm/rccl commit: 70804da15b]
This commit is contained in:
Wenkai Du
2019-07-05 14:29:20 -07:00
committed by GitHub
orang tua 20975921dd
melakukan b7322c800a
1 mengubah file dengan 258 tambahan dan 195 penghapusan
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl/tools/rccl-prim-test/rccl_prim_test.cpp
+258 -195
Melihat File
@@ -33,6 +33,7 @@ THE SOFTWARE.
#include <hip/hip_runtime.h>
#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<int op>
__global__ void flag_sync_kernel(struct transfer_data_t* transfer_data, struct profiling_data_t* profiling_data) {
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 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<COPY_UNROLL, THREADS, float>(transfer_data->dest0 + offset, transfer_data->src0 + offset, n);
if (op == OP_LOCALCOPY) Copy<COPY_UNROLL, THREADS, float>(transfer_data->dest1 + offset, transfer_data->src0 + offset, n);
if (op == OP_DOUBLECOPY) DoubleCopy<DOUBLECOPY_UNROLL, THREADS, float>(transfer_data->dest0 + offset, transfer_data->dest1 + offset, transfer_data->src0 + offset, n);
if (op == OP_REDUCE) Reduce<REDUCE_UNROLL, THREADS, float>(transfer_data->dest0 + offset, transfer_data->src0 + offset, transfer_data->src1 + offset, n);
if (op == OP_REDUCECOPY) ReduceCopy<REDUCECOPY_UNROLL, THREADS, float>(transfer_data->dest0 + offset, transfer_data->dest1 + offset, transfer_data->src0 + offset, transfer_data->src1 + offset, n);
if (op == OP_COPY) Copy<COPY_UNROLL, THREADS, float>(transfer_data->dest0[bid], transfer_data->src0[bid], n);
if (op == OP_LOCALCOPY) Copy<COPY_UNROLL, THREADS, float>(transfer_data->dest1[bid], transfer_data->src0[bid], n);
if (op == OP_DOUBLECOPY) DoubleCopy<DOUBLECOPY_UNROLL, THREADS, float>(transfer_data->dest0[bid], transfer_data->dest1[bid], transfer_data->src0[bid], n);
if (op == OP_REDUCE) Reduce<REDUCE_UNROLL, THREADS, float>(transfer_data->dest0[bid], transfer_data->src0[bid], transfer_data->src1[bid], n);
if (op == OP_REDUCECOPY) ReduceCopy<REDUCECOPY_UNROLL, THREADS, float>(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<OP_COPY>,
flag_sync_kernel<OP_LOCALCOPY>,
flag_sync_kernel<OP_DOUBLECOPY>,
flag_sync_kernel<OP_REDUCE>,
flag_sync_kernel<OP_REDUCECOPY>,
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_REDUCE, 0>,
flag_sync_kernel<OP_REDUCE, 1>,
flag_sync_kernel<OP_REDUCECOPY, 0>,
flag_sync_kernel<OP_REDUCECOPY, 1>,
};
__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<std::chrono::duration<double>>(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));
}