From 5a2f74b8d097629c8cd8b6724b9f056b542a535f Mon Sep 17 00:00:00 2001 From: rohit pathania Date: Fri, 9 Aug 2019 12:44:06 +0530 Subject: [PATCH 1/4] Adding linkinfo and srcGPU to destGPU info --- tools/rccl-prim-test/rccl_prim_test.cpp | 29 +++++++++++++++++++++---- 1 file changed, 25 insertions(+), 4 deletions(-) diff --git a/tools/rccl-prim-test/rccl_prim_test.cpp b/tools/rccl-prim-test/rccl_prim_test.cpp index b63fd058eb..3e449a2fd7 100644 --- a/tools/rccl-prim-test/rccl_prim_test.cpp +++ b/tools/rccl-prim-test/rccl_prim_test.cpp @@ -234,6 +234,21 @@ bool cmdOptionExists(char** begin, char** end, const std::string& option) { return std::find(begin, end, option) != end; } +std::string link_type(uint32_t type) { + switch(type) { + case HSA_AMD_LINK_INFO_TYPE_QPI : + return "QPI"; + case HSA_AMD_LINK_INFO_TYPE_INFINBAND: + return "INFINBAND"; + case HSA_AMD_LINK_INFO_TYPE_XGMI: + return "XGMI"; + case HSA_AMD_LINK_INFO_TYPE_PCIE: + return "PCIE"; + default: + return "HYPERTRANSPORT"; + } +} + int main(int argc,char* argv[]) { if (cmdOptionExists(argv, argv + argc, "-h")) { @@ -365,7 +380,7 @@ int main(int argc,char* argv[]) 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]); + printf("[Testing %s]: \n", OpsName[op]); // 2 warm up cycles for (int i = 0; i < 2; i ++) { for (int i = 0; i < nGpu; i ++) { @@ -410,17 +425,23 @@ int main(int argc,char* argv[]) auto delta = std::chrono::high_resolution_clock::now() - start; double deltaSec = std::chrono::duration_cast>(delta).count(); + std::cout<<"***GPU to GPU Transfer Profiling Data***"<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)); + fprintf(stderr, "[GPU %d -> GPU %d][%s]:time %.4fs bytes_transferred %lu kernel throughput %.2f GB/s\n", + i, next_gpu,link_type(linktype).c_str(),t0, profiling_data[i]->bytes_transferred, (double)profiling_data[i]->bytes_transferred/(t0*1.0E9)); } - + std::cout<<"***Application Level Transfer Profiling Data***"<bytes_transferred) / (deltaSec*1.0E9); printf("Transfered %lu bytes in %f s. Throughput %f GB/s\n", profiling_data[0]->bytes_transferred, deltaSec, speed); } From 3bbf924ff8c0948db21b2092ad69a62557a3461c Mon Sep 17 00:00:00 2001 From: rohit pathania Date: Fri, 9 Aug 2019 12:44:06 +0530 Subject: [PATCH 2/4] Adding linkinfo and srcGPU to destGPU info --- tools/rccl-prim-test/rccl_prim_test.cpp | 17 +++++++++++++---- 1 file changed, 13 insertions(+), 4 deletions(-) diff --git a/tools/rccl-prim-test/rccl_prim_test.cpp b/tools/rccl-prim-test/rccl_prim_test.cpp index b63fd058eb..8a43b05111 100644 --- a/tools/rccl-prim-test/rccl_prim_test.cpp +++ b/tools/rccl-prim-test/rccl_prim_test.cpp @@ -234,6 +234,9 @@ bool cmdOptionExists(char** begin, char** end, const std::string& option) { return std::find(begin, end, option) != end; } + +static const char* link_type_name[] = {"HT", "QPI", "PCIE", "IB", "XGMI"}; + int main(int argc,char* argv[]) { if (cmdOptionExists(argv, argv + argc, "-h")) { @@ -365,7 +368,7 @@ int main(int argc,char* argv[]) 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]); + printf("[Testing %s]: \n", OpsName[op]); // 2 warm up cycles for (int i = 0; i < 2; i ++) { for (int i = 0; i < nGpu; i ++) { @@ -410,17 +413,23 @@ int main(int argc,char* argv[]) auto delta = std::chrono::high_resolution_clock::now() - start; double deltaSec = std::chrono::duration_cast>(delta).count(); + std::cout<<"***GPU to GPU Transfer Profiling Data***"<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)); + fprintf(stderr, "[GPU %d -> GPU %d][%s]:time %.4fs bytes_transferred %lu kernel throughput %.2f GB/s\n", + i, next_gpu,link_type_name[linktype],t0, profiling_data[i]->bytes_transferred, (double)profiling_data[i]->bytes_transferred/(t0*1.0E9)); } - + std::cout<<"***Application Level Transfer Profiling Data***"<bytes_transferred) / (deltaSec*1.0E9); printf("Transfered %lu bytes in %f s. Throughput %f GB/s\n", profiling_data[0]->bytes_transferred, deltaSec, speed); } From 65e2f5d87b74771c66d7a095c22f0bf738236f8c Mon Sep 17 00:00:00 2001 From: rohit pathania Date: Wed, 14 Aug 2019 12:55:12 +0530 Subject: [PATCH 3/4] Modified the code to use RTC clock frequency based on gpu gcn id --- tools/rccl-prim-test/rccl_prim_test.cpp | 27 ++++++++++++++++++++----- 1 file changed, 22 insertions(+), 5 deletions(-) diff --git a/tools/rccl-prim-test/rccl_prim_test.cpp b/tools/rccl-prim-test/rccl_prim_test.cpp index ae4b768be1..57d3bc6bfd 100644 --- a/tools/rccl-prim-test/rccl_prim_test.cpp +++ b/tools/rccl-prim-test/rccl_prim_test.cpp @@ -42,6 +42,13 @@ THE SOFTWARE. #define DOUBLECOPY_UNROLL 2 #define REDUCECOPY_UNROLL 2 +#define RTC_CLOCK_FREQ_VEGA20 2.7E07 +//Right now kept the MI100 RTC frequency same as Vega20 +//as we are not aware of MI100 frequency, once we we come to know about it +//we will update it. +#define RTC_CLOCK_FREQ_MI100 2.7E07 +#define RTC_CLOCK_FREQ_DEFAULT 2.7E07 + struct transfer_data_t { float *dest0[MAX_WORKGROUPS]; //remote fine grain float *src0[MAX_WORKGROUPS]; //local fine grain @@ -420,15 +427,25 @@ int main(int argc,char* argv[]) sizeof(struct profiling_data_t), hipMemcpyDeviceToHost, stream[i])); HIPCHECK(hipStreamSynchronize(stream[i])); -#define RTC_CLOCK_FREQ 2.7E07 int next_gpu = findNextGpu(ring_0, i, nGpu); uint32_t linktype; uint32_t hopcount; HIPCHECK(hipExtGetLinkTypeAndHopCount(i, next_gpu , &linktype, &hopcount)); - - double t0 = (double)profiling_data[i]->write_cycles/((double)RTC_CLOCK_FREQ)/(double)workgroups; - fprintf(stderr, "[GPU %d -> GPU %d][%s]:time %.4fs bytes_transferred %lu kernel throughput %.2f GB/s\n", - i, next_gpu,link_type_name[linktype],t0, profiling_data[i]->bytes_transferred, (double)profiling_data[i]->bytes_transferred/(t0*1.0E9)); + hipDeviceProp_t prop; + HIPCHECK(hipGetDeviceProperties(&prop, i)); + if(prop.gcnArch == 906 ) { + double t0 = (double)profiling_data[i]->write_cycles/((double)RTC_CLOCK_FREQ_VEGA20)/(double)workgroups; + fprintf(stderr, "[GPU %d -> GPU %d][%s]:time %.4fs bytes_transferred %lu kernel throughput %.2f GB/s\n", + i, next_gpu,link_type_name[linktype],t0, profiling_data[i]->bytes_transferred, (double)profiling_data[i]->bytes_transferred/(t0*1.0E9)); + } else if (prop.gcnArch == 908 ){ + double t0 = (double)profiling_data[i]->write_cycles/((double)RTC_CLOCK_FREQ_MI100)/(double)workgroups; + fprintf(stderr, "[GPU %d -> GPU %d][%s]:time %.4fs bytes_transferred %lu kernel throughput %.2f GB/s\n", + i, next_gpu,link_type_name[linktype],t0, profiling_data[i]->bytes_transferred, (double)profiling_data[i]->bytes_transferred/(t0*1.0E9)); + } else { + double t0 = (double)profiling_data[i]->write_cycles/((double)RTC_CLOCK_FREQ_DEFAULT)/(double)workgroups; + fprintf(stderr, "[GPU %d -> GPU %d][%s]:time %.4fs bytes_transferred %lu kernel throughput %.2f GB/s\n", + i, next_gpu,link_type_name[linktype],t0, profiling_data[i]->bytes_transferred, (double)profiling_data[i]->bytes_transferred/(t0*1.0E9)); + } } std::cout<<"***Application Level Transfer Profiling Data***"<bytes_transferred) / (deltaSec*1.0E9); From 40e30b5168574238bb30bd4717e80d99d1619dd7 Mon Sep 17 00:00:00 2001 From: rpathani Date: Mon, 19 Aug 2019 12:44:11 +0530 Subject: [PATCH 4/4] Update rccl_prim_test.cpp --- tools/rccl-prim-test/rccl_prim_test.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tools/rccl-prim-test/rccl_prim_test.cpp b/tools/rccl-prim-test/rccl_prim_test.cpp index 0e0fe3c63c..d57700474b 100644 --- a/tools/rccl-prim-test/rccl_prim_test.cpp +++ b/tools/rccl-prim-test/rccl_prim_test.cpp @@ -43,7 +43,7 @@ THE SOFTWARE. #define REDUCECOPY_UNROLL 2 #define RTC_CLOCK_FREQ_VEGA20 2.7E07 -//Right now kept the MI100 RTC frequency same as Vega20 +//Right now kept the MI100 RTC frequency same as Vega20 //as we are not aware of MI100 frequency, once we we come to know about it //we will update it. #define RTC_CLOCK_FREQ_MI100 2.7E07