diff --git a/projects/rccl/tools/p2p-latency-test/build_and_run.sh b/projects/rccl/tools/p2p-latency-test/build_and_run.sh index e2aec1719f..c3e87d202b 100644 --- a/projects/rccl/tools/p2p-latency-test/build_and_run.sh +++ b/projects/rccl/tools/p2p-latency-test/build_and_run.sh @@ -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 diff --git a/projects/rccl/tools/p2p-latency-test/ll_latency_test.cpp b/projects/rccl/tools/p2p-latency-test/ll_latency_test.cpp index 17a952eb81..fb10f3fe66 100644 --- a/projects/rccl/tools/p2p-latency-test/ll_latency_test.cpp +++ b/projects/rccl/tools/p2p-latency-test/ll_latency_test.cpp @@ -11,6 +11,8 @@ #include #include #include +#include //cerr +#include #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=%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; -} +} \ No newline at end of file diff --git a/projects/rccl/tools/p2p-latency-test/p2p_latency_test.cpp b/projects/rccl/tools/p2p-latency-test/p2p_latency_test.cpp index b9d304f02c..7f1e6bc58c 100644 --- a/projects/rccl/tools/p2p-latency-test/p2p_latency_test.cpp +++ b/projects/rccl/tools/p2p-latency-test/p2p_latency_test.cpp @@ -10,6 +10,8 @@ #include #include #include +#include //cerr +#include #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=%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; }