p2p/ll-latency-test: convert to single thread tests (#857)
[ROCm/rccl commit: 148e3430f4]
This commit is contained in:
Wenkai Du
GitHub
@@ -4,21 +4,22 @@
|
||||
make
|
||||
|
||||
# Example run: test one-way latency between GPU 0 and GPU 1 in both directions.
|
||||
export HSA_FORCE_FINE_GRAIN_PCIE=1
|
||||
|
||||
echo Running p2p_latency_test using GPU pair 0 1
|
||||
HIP_VISIBLE_DEVICES=0 ./p2p_latency_test 0 & HIP_VISIBLE_DEVICES=1 ./p2p_latency_test 1
|
||||
./p2p_latency_test 0 1
|
||||
|
||||
sleep 1
|
||||
|
||||
echo Running p2p_latency_test using GPU pair 1 0
|
||||
HIP_VISIBLE_DEVICES=1 ./p2p_latency_test 0 & HIP_VISIBLE_DEVICES=0 ./p2p_latency_test 1
|
||||
./p2p_latency_test 1 0
|
||||
|
||||
sleep 1
|
||||
|
||||
echo Running ll_latency_test using GPU pair 0 1
|
||||
HIP_VISIBLE_DEVICES=0 ./ll_latency_test 0 & HIP_VISIBLE_DEVICES=1 ./ll_latency_test 1
|
||||
./ll_latency_test 0 1
|
||||
|
||||
sleep 1
|
||||
|
||||
echo Running ll_latency_test using GPU pair 1 0
|
||||
HIP_VISIBLE_DEVICES=1 ./ll_latency_test 0 & HIP_VISIBLE_DEVICES=0 ./ll_latency_test 1
|
||||
./ll_latency_test 1 0
|
||||
|
||||
@@ -11,6 +11,8 @@
|
||||
#include <string>
|
||||
#include <unistd.h>
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <iostream> //cerr
|
||||
#include <cstring>
|
||||
|
||||
#define HIP_IPC_MEM_MIN_SIZE 2097152UL
|
||||
|
||||
@@ -49,9 +51,7 @@ __device__ void storeLL(union LLFifoLine* dst, uint64_t val, uint32_t flag) {
|
||||
|
||||
#define LL_SPINS_BEFORE_CHECK_ABORT 1000000
|
||||
|
||||
__device__ uint32_t *abortFlag;
|
||||
|
||||
inline __device__ int checkAbort(int &spins) {
|
||||
inline __device__ int checkAbort(int &spins, uint32_t* abortFlag) {
|
||||
uint32_t abort = 0;
|
||||
spins++;
|
||||
if (spins == LL_SPINS_BEFORE_CHECK_ABORT) {
|
||||
@@ -61,7 +61,7 @@ inline __device__ int checkAbort(int &spins) {
|
||||
return abort;
|
||||
}
|
||||
|
||||
__device__ uint64_t readLL(union LLFifoLine* src, uint32_t flag) {
|
||||
__device__ uint64_t readLL(union LLFifoLine* src, uint32_t flag, uint32_t* abortFlag) {
|
||||
uint32_t data1, flag1, data2, flag2;
|
||||
int spins = 0;
|
||||
|
||||
@@ -69,7 +69,7 @@ __device__ uint64_t readLL(union LLFifoLine* src, uint32_t flag) {
|
||||
do {
|
||||
i4.v[0] = __builtin_nontemporal_load(src->v);
|
||||
i4.v[1] = __builtin_nontemporal_load(src->v+1);
|
||||
if (checkAbort(spins)) break;
|
||||
if (checkAbort(spins, abortFlag)) break;
|
||||
} while ((i4.flag1 != flag) || (i4.flag2 != flag));
|
||||
uint64_t val64 = (uint64_t)(i4.data1) + (((uint64_t)i4.data2) << 32);
|
||||
|
||||
@@ -77,37 +77,37 @@ __device__ uint64_t readLL(union LLFifoLine* src, uint32_t flag) {
|
||||
}
|
||||
|
||||
|
||||
__global__ void PingKernel(LLFifoLine* local_flag, LLFifoLine* remote_flag, uint64_t* time_delta) {
|
||||
__global__ void PingKernel(LLFifoLine* local_flag, LLFifoLine* remote_flag, uint64_t* time_delta, uint32_t* abortFlag) {
|
||||
int tid = threadIdx.x;
|
||||
#pragma unroll
|
||||
for (uint32_t i = 1; i < NUM_LOOPS_WARMUP; i++) {
|
||||
storeLL(remote_flag+tid, i, i);
|
||||
while (readLL(local_flag+tid, i) != i);
|
||||
while (readLL(local_flag+tid, i, abortFlag) != i);
|
||||
}
|
||||
uint64_t start_time, end_time;
|
||||
if (tid == 0) start_time = wall_clock64();
|
||||
#pragma unroll
|
||||
for (uint32_t i = NUM_LOOPS_WARMUP; i <= NUM_LOOPS_WARMUP + NUM_LOOPS_RUN; i++) {
|
||||
storeLL(remote_flag+tid, i, i);
|
||||
while (readLL(local_flag+tid, i) != i);
|
||||
while (readLL(local_flag+tid, i, abortFlag) != i);
|
||||
}
|
||||
__syncthreads();
|
||||
if (tid == 0) end_time = wall_clock64();
|
||||
if (tid == 0) *time_delta = end_time - start_time;
|
||||
}
|
||||
|
||||
__global__ void PongKernel(LLFifoLine* local_flag, LLFifoLine* remote_flag, uint64_t* time_delta) {
|
||||
__global__ void PongKernel(LLFifoLine* local_flag, LLFifoLine* remote_flag, uint64_t* time_delta, uint32_t* abortFlag) {
|
||||
int tid = threadIdx.x;
|
||||
#pragma unroll
|
||||
for (uint32_t i = 1; i < NUM_LOOPS_WARMUP; i++) {
|
||||
while (readLL(local_flag+tid, i) != i);
|
||||
while (readLL(local_flag+tid, i, abortFlag) != i);
|
||||
storeLL(remote_flag+tid, i, i);
|
||||
}
|
||||
uint64_t start_time, end_time;
|
||||
if (tid == 0) start_time = wall_clock64();
|
||||
#pragma unroll
|
||||
for (uint32_t i = NUM_LOOPS_WARMUP; i <= NUM_LOOPS_WARMUP + NUM_LOOPS_RUN; i++) {
|
||||
while (readLL(local_flag+tid, i) != i);
|
||||
while (readLL(local_flag+tid, i, abortFlag) != i);
|
||||
storeLL(remote_flag+tid, i, i);
|
||||
}
|
||||
__syncthreads();
|
||||
@@ -115,70 +115,79 @@ __global__ void PongKernel(LLFifoLine* local_flag, LLFifoLine* remote_flag, uint
|
||||
if (tid == 0) *time_delta = end_time - start_time;
|
||||
}
|
||||
|
||||
#define HIPCHECK(cmd) \
|
||||
do { \
|
||||
hipError_t error = (cmd); \
|
||||
if (error != hipSuccess) \
|
||||
{ \
|
||||
std::cerr << "Encountered HIP error (" << error << ") at line " \
|
||||
<< __LINE__ << " in file " << __FILE__ << "\n"; \
|
||||
exit(-1); \
|
||||
} \
|
||||
} while (0)
|
||||
|
||||
int main(int argc, char** argv) {
|
||||
hipStream_t stream;
|
||||
hipStream_t stream[2];
|
||||
hipError_t err = hipSuccess;
|
||||
int device_id[2];
|
||||
hipDeviceProp_t prop[2];
|
||||
|
||||
if (argc != 2) {
|
||||
fprintf(stderr, "Usage: ./p2p_latency_test <flag>; flag=%d for ping mode, flag=%d for pong mode\n", PING_MODE, PONG_MODE);
|
||||
if (argc != 3) {
|
||||
fprintf(stderr, "Usage: ./ll_latency_test ping_dev_id pong_dev_id\n");
|
||||
return -1;
|
||||
}
|
||||
device_id[0] = atoi(argv[1]);
|
||||
device_id[1] = atoi(argv[2]);
|
||||
|
||||
hipDeviceProp_t prop;
|
||||
int device_id;
|
||||
hipGetDevice(&device_id);
|
||||
hipGetDeviceProperties(&prop, device_id);
|
||||
fprintf(stdout, "Using devices %d %d\n", device_id[0], device_id[1]);
|
||||
|
||||
LLFifoLine *local_flag = nullptr;
|
||||
LLFifoLine *remote_flag = nullptr;
|
||||
uint64_t *time_delta = nullptr;
|
||||
hipStreamCreateWithFlags(&stream, hipStreamNonBlocking);
|
||||
hipExtMallocWithFlags((void**)&local_flag, HIP_IPC_MEM_MIN_SIZE, prop.gcnArch / 10 == 94 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained);
|
||||
hipExtMallocWithFlags((void**)&time_delta, HIP_IPC_MEM_MIN_SIZE, prop.gcnArch / 10 == 94 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained);
|
||||
hipMemsetAsync(local_flag, 0, HIP_IPC_MEM_MIN_SIZE, stream);
|
||||
hipMemsetAsync(time_delta, 0, HIP_IPC_MEM_MIN_SIZE, stream);
|
||||
hipStreamSynchronize(stream);
|
||||
hipIpcMemHandle_t local_handle;
|
||||
hipIpcMemHandle_t remote_handle;
|
||||
hipIpcGetMemHandle(&local_handle, local_flag);
|
||||
LLFifoLine *flag[2];
|
||||
uint64_t *time_delta[2];
|
||||
uint32_t *abortFlag[2];
|
||||
|
||||
const char* ping_file_path = "/tmp/ping_ipc_handle";
|
||||
const char* pong_file_path = "/tmp/pong_ipc_handle";
|
||||
const char* file_paths[2] = {ping_file_path, pong_file_path};
|
||||
int self_mode = atoi(argv[1]);
|
||||
if (self_mode == PING_MODE || self_mode == PONG_MODE) {
|
||||
int peer_mode = 1 - self_mode;
|
||||
auto self_file = std::fstream(file_paths[self_mode], std::ios::out | std::ios::binary);
|
||||
self_file.write((char*)&local_handle, sizeof(local_handle));
|
||||
self_file.close();
|
||||
sleep(5);
|
||||
auto peer_file = std::fstream(file_paths[peer_mode], std::ios::in | std::ios::binary);
|
||||
peer_file.read((char*)&remote_handle, sizeof(remote_handle));
|
||||
peer_file.close();
|
||||
err = hipIpcOpenMemHandle((void**)(&remote_flag), remote_handle, hipIpcMemLazyEnablePeerAccess);
|
||||
if (err != hipSuccess) {
|
||||
fprintf(stderr, "hipIpcOpenMemHandle error %d\n", (int)err);
|
||||
return -1;
|
||||
}
|
||||
if (self_mode == PING_MODE) {
|
||||
PingKernel<<<1, LL_MAX_THREAD, 0, stream>>>(local_flag, remote_flag, time_delta);
|
||||
} else {
|
||||
PongKernel<<<1, LL_MAX_THREAD, 0, stream>>>(local_flag, remote_flag, time_delta);
|
||||
}
|
||||
err = hipStreamSynchronize(stream);
|
||||
if (err != hipSuccess) {
|
||||
fprintf(stderr, "hipStreamSynchronize error %d\n", (int)err);
|
||||
return -1;
|
||||
}
|
||||
if (self_mode == PING_MODE) {
|
||||
double vega_gpu_rtc_freq = (prop.gcnArch / 10 == 94) ? 1.0E8 : 2.5E7;
|
||||
fprintf(stdout, "One-way latency in us: %g\n", double(*time_delta) * 1e6 / NUM_LOOPS_RUN / vega_gpu_rtc_freq / 2);
|
||||
}
|
||||
std::remove(file_paths[self_mode]);
|
||||
} else {
|
||||
fprintf(stderr, "Invalid mode %d\n", self_mode);
|
||||
return -1;
|
||||
}
|
||||
HIPCHECK(hipSetDevice(device_id[0]));
|
||||
HIPCHECK(hipStreamCreateWithFlags(&stream[0], hipStreamNonBlocking));
|
||||
HIPCHECK(hipDeviceEnablePeerAccess(device_id[1], 0));
|
||||
HIPCHECK(hipGetDeviceProperties(&prop[0], device_id[0]));
|
||||
HIPCHECK(hipExtMallocWithFlags((void**)&flag[0], HIP_IPC_MEM_MIN_SIZE, prop[0].gcnArch / 10 == 94 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
|
||||
HIPCHECK(hipMalloc((void**)&time_delta[0], HIP_IPC_MEM_MIN_SIZE));
|
||||
HIPCHECK(hipMalloc((void**)&abortFlag[0], sizeof(uint32_t)));
|
||||
HIPCHECK(hipMemsetAsync(flag[0], 0, HIP_IPC_MEM_MIN_SIZE, stream[0]));
|
||||
HIPCHECK(hipMemsetAsync(abortFlag[0], 0, sizeof(uint32_t), stream[0]));
|
||||
HIPCHECK(hipStreamSynchronize(stream[0]));
|
||||
|
||||
HIPCHECK(hipSetDevice(device_id[1]));
|
||||
HIPCHECK(hipStreamCreateWithFlags(&stream[1], hipStreamNonBlocking));
|
||||
HIPCHECK(hipDeviceEnablePeerAccess(device_id[0], 0));
|
||||
HIPCHECK(hipGetDeviceProperties(&prop[1], device_id[1]));
|
||||
HIPCHECK(hipExtMallocWithFlags((void**)&flag[1], HIP_IPC_MEM_MIN_SIZE, prop[1].gcnArch / 10 == 94 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
|
||||
HIPCHECK(hipMalloc((void**)&time_delta[1], HIP_IPC_MEM_MIN_SIZE));
|
||||
HIPCHECK(hipMalloc((void**)&abortFlag[1], sizeof(uint32_t)));
|
||||
HIPCHECK(hipMemsetAsync(flag[1], 0, HIP_IPC_MEM_MIN_SIZE, stream[1]));
|
||||
HIPCHECK(hipMemsetAsync(abortFlag[1], 0, sizeof(uint32_t), stream[0]));
|
||||
HIPCHECK(hipStreamSynchronize(stream[1]));
|
||||
|
||||
HIPCHECK(hipSetDevice(device_id[0]));
|
||||
PingKernel<<<1, LL_MAX_THREAD, 0, stream[0]>>>(flag[0], flag[1], time_delta[0], abortFlag[0]);
|
||||
|
||||
HIPCHECK(hipSetDevice(device_id[1]));
|
||||
PongKernel<<<1, LL_MAX_THREAD, 0, stream[1]>>>(flag[1], flag[0], time_delta[1], abortFlag[1]);
|
||||
|
||||
double vega_gpu_rtc_freq;
|
||||
|
||||
HIPCHECK(hipStreamSynchronize(stream[0]));
|
||||
vega_gpu_rtc_freq = (prop[0].gcnArch / 10 == 94) ? 1.0E8 : 2.5E7;
|
||||
fprintf(stdout, "One-way latency in us: %g\n", double(*time_delta[0]) * 1e6 / NUM_LOOPS_RUN / vega_gpu_rtc_freq / 2);
|
||||
|
||||
HIPCHECK(hipStreamSynchronize(stream[1]));
|
||||
vega_gpu_rtc_freq = (prop[1].gcnArch / 10 == 94) ? 1.0E8 : 2.5E7;
|
||||
fprintf(stdout, "One-way latency in us: %g\n", double(*time_delta[1]) * 1e6 / NUM_LOOPS_RUN / vega_gpu_rtc_freq / 2);
|
||||
|
||||
HIPCHECK(hipFree(flag[0]));
|
||||
HIPCHECK(hipFree(time_delta[0]));
|
||||
HIPCHECK(hipFree(abortFlag[0]));
|
||||
HIPCHECK(hipFree(flag[1]));
|
||||
HIPCHECK(hipFree(time_delta[1]));
|
||||
HIPCHECK(hipFree(abortFlag[1]));
|
||||
return 0;
|
||||
}
|
||||
}
|
||||
@@ -10,6 +10,8 @@
|
||||
#include <string>
|
||||
#include <unistd.h>
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <iostream> //cerr
|
||||
#include <cstring>
|
||||
|
||||
#define HIP_IPC_MEM_MIN_SIZE 2097152UL
|
||||
|
||||
@@ -51,70 +53,72 @@ __global__ void PongKernel(uint64_t* local_flag, uint64_t* remote_flag, uint64_t
|
||||
*time_delta = end_time - start_time;
|
||||
}
|
||||
|
||||
#define HIPCHECK(cmd) \
|
||||
do { \
|
||||
hipError_t error = (cmd); \
|
||||
if (error != hipSuccess) \
|
||||
{ \
|
||||
std::cerr << "Encountered HIP error (" << error << ") at line " \
|
||||
<< __LINE__ << " in file " << __FILE__ << "\n"; \
|
||||
exit(-1); \
|
||||
} \
|
||||
} while (0)
|
||||
|
||||
int main(int argc, char** argv) {
|
||||
hipStream_t stream;
|
||||
hipStream_t stream[2];
|
||||
hipError_t err = hipSuccess;
|
||||
int device_id[2];
|
||||
hipDeviceProp_t prop[2];
|
||||
|
||||
if (argc != 2) {
|
||||
fprintf(stderr, "Usage: ./ll_latency_test <flag>; flag=%d for ping mode, flag=%d for pong mode\n", PING_MODE, PONG_MODE);
|
||||
if (argc != 3) {
|
||||
fprintf(stderr, "Usage: ./p2p_latency_test ping_dev_id pong_dev_id\n");
|
||||
return -1;
|
||||
}
|
||||
device_id[0] = atoi(argv[1]);
|
||||
device_id[1] = atoi(argv[2]);
|
||||
|
||||
hipDeviceProp_t prop;
|
||||
int device_id;
|
||||
hipGetDevice(&device_id);
|
||||
hipGetDeviceProperties(&prop, device_id);
|
||||
fprintf(stdout, "Using devices %d %d\n", device_id[0], device_id[1]);
|
||||
|
||||
uint64_t *local_flag = nullptr;
|
||||
uint64_t *remote_flag = nullptr;
|
||||
uint64_t *time_delta = nullptr;
|
||||
hipStreamCreateWithFlags(&stream, hipStreamNonBlocking);
|
||||
hipExtMallocWithFlags((void**)&local_flag, HIP_IPC_MEM_MIN_SIZE, prop.gcnArch / 10 == 94 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained);
|
||||
hipExtMallocWithFlags((void**)&time_delta, HIP_IPC_MEM_MIN_SIZE, prop.gcnArch / 10 == 94 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained);
|
||||
hipMemsetAsync(local_flag, 0, HIP_IPC_MEM_MIN_SIZE, stream);
|
||||
hipMemsetAsync(time_delta, 0, HIP_IPC_MEM_MIN_SIZE, stream);
|
||||
hipStreamSynchronize(stream);
|
||||
hipIpcMemHandle_t local_handle;
|
||||
hipIpcMemHandle_t remote_handle;
|
||||
hipIpcGetMemHandle(&local_handle, local_flag);
|
||||
uint64_t *flag[2];
|
||||
uint64_t *time_delta[2];
|
||||
|
||||
const char* ping_file_path = "/tmp/ping_ipc_handle";
|
||||
const char* pong_file_path = "/tmp/pong_ipc_handle";
|
||||
const char* file_paths[2] = {ping_file_path, pong_file_path};
|
||||
int self_mode = atoi(argv[1]);
|
||||
if (self_mode == PING_MODE || self_mode == PONG_MODE) {
|
||||
int peer_mode = 1 - self_mode;
|
||||
auto self_file = std::fstream(file_paths[self_mode], std::ios::out | std::ios::binary);
|
||||
self_file.write((char*)&local_handle, sizeof(local_handle));
|
||||
self_file.close();
|
||||
sleep(5);
|
||||
auto peer_file = std::fstream(file_paths[peer_mode], std::ios::in | std::ios::binary);
|
||||
peer_file.read((char*)&remote_handle, sizeof(remote_handle));
|
||||
peer_file.close();
|
||||
err = hipIpcOpenMemHandle((void**)(&remote_flag), remote_handle, hipIpcMemLazyEnablePeerAccess);
|
||||
if (err != hipSuccess) {
|
||||
fprintf(stderr, "hipIpcOpenMemHandle error %d\n", (int)err);
|
||||
return -1;
|
||||
}
|
||||
if (self_mode == PING_MODE) {
|
||||
PingKernel<<<1, 1, 0, stream>>>(local_flag, remote_flag, time_delta);
|
||||
} else {
|
||||
PongKernel<<<1, 1, 0, stream>>>(local_flag, remote_flag, time_delta);
|
||||
}
|
||||
err = hipStreamSynchronize(stream);
|
||||
if (err != hipSuccess) {
|
||||
fprintf(stderr, "hipStreamSynchronize error %d\n", (int)err);
|
||||
return -1;
|
||||
}
|
||||
if (self_mode == PING_MODE) {
|
||||
double vega_gpu_rtc_freq = (prop.gcnArch / 10 == 94) ? 1.0E8 : 2.5E7;
|
||||
fprintf(stdout, "One-way latency in us: %g\n", double(*time_delta) * 1e6 / NUM_LOOPS_RUN / vega_gpu_rtc_freq / 2);
|
||||
}
|
||||
std::remove(file_paths[self_mode]);
|
||||
} else {
|
||||
fprintf(stderr, "Invalid mode %d\n", self_mode);
|
||||
return -1;
|
||||
}
|
||||
HIPCHECK(hipSetDevice(device_id[0]));
|
||||
HIPCHECK(hipStreamCreateWithFlags(&stream[0], hipStreamNonBlocking));
|
||||
HIPCHECK(hipDeviceEnablePeerAccess(device_id[1], 0));
|
||||
HIPCHECK(hipGetDeviceProperties(&prop[0], device_id[0]));
|
||||
HIPCHECK(hipExtMallocWithFlags((void**)&flag[0], HIP_IPC_MEM_MIN_SIZE, prop[0].gcnArch / 10 == 94 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
|
||||
HIPCHECK(hipMalloc((void**)&time_delta[0], HIP_IPC_MEM_MIN_SIZE));
|
||||
HIPCHECK(hipMemsetAsync(flag[0], 0, HIP_IPC_MEM_MIN_SIZE, stream[0]));
|
||||
HIPCHECK(hipStreamSynchronize(stream[0]));
|
||||
|
||||
HIPCHECK(hipSetDevice(device_id[1]));
|
||||
HIPCHECK(hipStreamCreateWithFlags(&stream[1], hipStreamNonBlocking));
|
||||
HIPCHECK(hipDeviceEnablePeerAccess(device_id[0], 0));
|
||||
HIPCHECK(hipGetDeviceProperties(&prop[1], device_id[1]));
|
||||
HIPCHECK(hipExtMallocWithFlags((void**)&flag[1], HIP_IPC_MEM_MIN_SIZE, prop[1].gcnArch / 10 == 94 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
|
||||
HIPCHECK(hipMalloc((void**)&time_delta[1], HIP_IPC_MEM_MIN_SIZE));
|
||||
HIPCHECK(hipMemsetAsync(flag[1], 0, HIP_IPC_MEM_MIN_SIZE, stream[1]));
|
||||
HIPCHECK(hipStreamSynchronize(stream[1]));
|
||||
|
||||
HIPCHECK(hipSetDevice(device_id[0]));
|
||||
PingKernel<<<1, 1, 0, stream[0]>>>(flag[0], flag[1], time_delta[0]);
|
||||
|
||||
HIPCHECK(hipSetDevice(device_id[1]));
|
||||
PongKernel<<<1, 1, 0, stream[1]>>>(flag[1], flag[0], time_delta[1]);
|
||||
|
||||
double vega_gpu_rtc_freq;
|
||||
|
||||
HIPCHECK(hipStreamSynchronize(stream[0]));
|
||||
vega_gpu_rtc_freq = (prop[0].gcnArch / 10 == 94) ? 1.0E8 : 2.5E7;
|
||||
fprintf(stdout, "One-way latency in us: %g\n", double(*time_delta[0]) * 1e6 / NUM_LOOPS_RUN / vega_gpu_rtc_freq / 2);
|
||||
|
||||
HIPCHECK(hipStreamSynchronize(stream[1]));
|
||||
vega_gpu_rtc_freq = (prop[1].gcnArch / 10 == 94) ? 1.0E8 : 2.5E7;
|
||||
fprintf(stdout, "One-way latency in us: %g\n", double(*time_delta[1]) * 1e6 / NUM_LOOPS_RUN / vega_gpu_rtc_freq / 2);
|
||||
|
||||
HIPCHECK(hipFree(flag[0]));
|
||||
HIPCHECK(hipFree(time_delta[0]));
|
||||
HIPCHECK(hipFree(flag[1]));
|
||||
HIPCHECK(hipFree(time_delta[1]));
|
||||
return 0;
|
||||
}
|
||||
|
||||
Reference in New Issue
Block a user