From e1a3c5a6dc914789750bb55b90f3d83934c7d774 Mon Sep 17 00:00:00 2001 From: saurabhAMD Date: Tue, 7 May 2024 11:09:32 -0500 Subject: [PATCH 1/3] Cache flush [ROCm/rccl-tests commit: 3c0728e8ebecac1f2935806404da3b32bf083910] --- projects/rccl-tests/src/common.cu | 53 ++++++++++++++++++++++++++++++- projects/rccl-tests/src/common.h | 1 + 2 files changed, 53 insertions(+), 1 deletion(-) diff --git a/projects/rccl-tests/src/common.cu b/projects/rccl-tests/src/common.cu index 691c2402dd..19c60913d8 100644 --- a/projects/rccl-tests/src/common.cu +++ b/projects/rccl-tests/src/common.cu @@ -103,9 +103,47 @@ static int average = 1; static int numDevices = 1; static int delay_inout_place = 0; static int enable_out_of_place = 1; +static int enable_cache_flush = 0; #define NUM_BLOCKS 32 +#ifndef CHECK_HIP_ERROR +#define CHECK_HIP_ERROR(error) \ + if(error != hipSuccess) \ + { \ + fprintf(stderr, \ + "Hip error: '%s'(%d) at %s:%d\n", \ + hipGetErrorString(error), \ + error, \ + __FILE__, \ + __LINE__); \ + exit(EXIT_FAILURE); \ + } +#endif + +extern "C" __global__ void flush_icache() +{ + printf("flush_icache called \n"); + asm __volatile__("s_icache_inv \n\t" + "s_nop 0 \n\t" + "s_nop 0 \n\t" + "s_nop 0 \n\t" + "s_nop 0 \n\t" + "s_nop 0 \n\t" + "s_nop 0 \n\t" + "s_nop 0 \n\t" + "s_nop 0 \n\t" + "s_nop 0 \n\t" + "s_nop 0 \n\t" + "s_nop 0 \n\t" + "s_nop 0 \n\t" + "s_nop 0 \n\t" + "s_nop 0 \n\t" + "s_nop 0 \n\t" + "s_nop 0 \n\t" :: + :); +} + static double parsesize(const char *value) { long long int units; double size; @@ -437,6 +475,13 @@ testResult_t startColl(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t } #endif + if(enable_cache_flush > 0 && (enable_cache_flush==1 || ((iter % enable_cache_flush) == 0))) { + hipDeviceProp_t deviceProps; + CHECK_HIP_ERROR(hipGetDeviceProperties(&deviceProps, 0)); + int32_t gpu_block3 = deviceProps.multiProcessorCount * 60; + hipLaunchKernelGGL(flush_icache, dim3(gpu_block3), dim3(64), 0, args->streams[i]); + } + TESTCHECK(args->collTest->runColl( (void*)(in_place ? recvBuff + args->sendInplaceOffset*rank : sendBuff), (void*)(in_place ? recvBuff + args->recvInplaceOffset*rank : recvBuff), @@ -896,6 +941,7 @@ int main(int argc, char* argv[]) { {"report_cputime", required_argument, 0, 'C'}, {"average", required_argument, 0, 'a'}, {"out_of_place", required_argument, 0, 'O'}, + {"cache_flush", required_argument, 0, 'F'}, {"help", no_argument, 0, 'h'}, {} }; @@ -903,7 +949,7 @@ int main(int argc, char* argv[]) { while(1) { int c; - c = getopt_long(argc, argv, "t:g:b:e:i:f:n:m:w:p:c:o:d:r:z:Y:T:G:C:O:a:y:s:u:h:q:", longopts, &longindex); + c = getopt_long(argc, argv, "t:g:b:e:i:f:n:m:w:p:c:o:d:r:z:Y:T:G:C:O:F:a:y:s:u:h:q:", longopts, &longindex); if (c == -1) break; @@ -1003,6 +1049,9 @@ int main(int argc, char* argv[]) { case 'O': enable_out_of_place = strtol(optarg, NULL, 0); break; + case 'F': + enable_cache_flush = strtol(optarg, NULL, 0); + break; case 'a': average = (int)strtol(optarg, NULL, 0); break; @@ -1042,6 +1091,7 @@ int main(int argc, char* argv[]) { "[-G,--cudagraph ] \n\t" "[-C,--report_cputime <0/1>] \n\t" "[-O,--out_of_place <0/1>] \n\t" + "[-F,--cache_flush <0/1>] \n\t" "[-a,--average <0/1/2/3> report average iteration time <0=RANK0/1=AVG/2=MIN/3=MAX>] \n\t" "[-q,--delay ] \n\t" "[-h,--help]\n", @@ -1252,6 +1302,7 @@ testResult_t run() { threads[t].args.comms=comms+t*nGpus; threads[t].args.streams=streams+t*nGpus; threads[t].args.enable_out_of_place=enable_out_of_place; + threads[t].args.enable_cache_flush = enable_cache_flush; threads[t].args.errors=errors+t; threads[t].args.bw=bw+t; threads[t].args.bw_count=bw_count+t; diff --git a/projects/rccl-tests/src/common.h b/projects/rccl-tests/src/common.h index 23dccebc7d..dc4aa4a7a2 100644 --- a/projects/rccl-tests/src/common.h +++ b/projects/rccl-tests/src/common.h @@ -127,6 +127,7 @@ struct threadArgs { int* gpus; int localRank; int enable_out_of_place; + int enable_cache_flush; void** sendbuffs; size_t sendBytes; size_t sendInplaceOffset; From ce8e61cc3b1269171a2352621d3e971420aa2f2d Mon Sep 17 00:00:00 2001 From: saurabhAMD Date: Tue, 7 May 2024 11:32:30 -0500 Subject: [PATCH 2/3] Enable cache flush after every -F iteration. Default : 0 (No cache flush) [ROCm/rccl-tests commit: 699478dadf35a3a7644704efadf2198f503ed8d5] --- projects/rccl-tests/README.md | 1 + projects/rccl-tests/src/common.cu | 1 - 2 files changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/rccl-tests/README.md b/projects/rccl-tests/README.md index 1eadbe768e..9bc34de49e 100644 --- a/projects/rccl-tests/README.md +++ b/projects/rccl-tests/README.md @@ -91,6 +91,7 @@ All tests support the same set of arguments : * `-c,--check ` perform count iterations, checking correctness of results on each iteration. This can be quite slow on large numbers of GPUs. Default : 1. * `-z,--blocking <0/1>` Make NCCL collective blocking, i.e. have CPUs wait and sync after each collective. Default : 0. * `-G,--cudagraph ` Capture iterations as a CUDA graph and then replay specified number of times. Default : 0. + * `-F,--cache_flush ` Enable cache flush after every -F iteration. Default : 0 (No cache flush). ## Unit tests diff --git a/projects/rccl-tests/src/common.cu b/projects/rccl-tests/src/common.cu index 19c60913d8..4dcac1cbd9 100644 --- a/projects/rccl-tests/src/common.cu +++ b/projects/rccl-tests/src/common.cu @@ -123,7 +123,6 @@ static int enable_cache_flush = 0; extern "C" __global__ void flush_icache() { - printf("flush_icache called \n"); asm __volatile__("s_icache_inv \n\t" "s_nop 0 \n\t" "s_nop 0 \n\t" From 5700751b7edd3fecb4d66e974ec62c9ef8a3692c Mon Sep 17 00:00:00 2001 From: saurabhAMD Date: Fri, 10 May 2024 08:46:13 -0700 Subject: [PATCH 3/3] updating cache flush on functionality [ROCm/rccl-tests commit: 74c4177f58113a61638bccd44df1b7d2d7915923] --- projects/rccl-tests/src/common.cu | 13 ++++++++----- 1 file changed, 8 insertions(+), 5 deletions(-) diff --git a/projects/rccl-tests/src/common.cu b/projects/rccl-tests/src/common.cu index 4dcac1cbd9..e545870545 100644 --- a/projects/rccl-tests/src/common.cu +++ b/projects/rccl-tests/src/common.cu @@ -24,6 +24,7 @@ #include "git_version.h" int test_ncclVersion = 0; // init'd with ncclGetVersion() +int32_t gpu_block3; #if NCCL_MAJOR >= 2 ncclDataType_t test_types[ncclNumTypes] = { @@ -474,10 +475,7 @@ testResult_t startColl(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t } #endif - if(enable_cache_flush > 0 && (enable_cache_flush==1 || ((iter % enable_cache_flush) == 0))) { - hipDeviceProp_t deviceProps; - CHECK_HIP_ERROR(hipGetDeviceProperties(&deviceProps, 0)); - int32_t gpu_block3 = deviceProps.multiProcessorCount * 60; + if(enable_cache_flush > 0 && ((iter % enable_cache_flush) == 0)) { hipLaunchKernelGGL(flush_icache, dim3(gpu_block3), dim3(64), 0, args->streams[i]); } @@ -1050,6 +1048,11 @@ int main(int argc, char* argv[]) { break; case 'F': enable_cache_flush = strtol(optarg, NULL, 0); + if (enable_cache_flush > 0) { + hipDeviceProp_t deviceProps; + CHECK_HIP_ERROR(hipGetDeviceProperties(&deviceProps, 0)); + gpu_block3 = deviceProps.multiProcessorCount * 60; + } break; case 'a': average = (int)strtol(optarg, NULL, 0); @@ -1090,7 +1093,7 @@ int main(int argc, char* argv[]) { "[-G,--cudagraph ] \n\t" "[-C,--report_cputime <0/1>] \n\t" "[-O,--out_of_place <0/1>] \n\t" - "[-F,--cache_flush <0/1>] \n\t" + "[-F,--cache_flush ] \n\t" "[-a,--average <0/1/2/3> report average iteration time <0=RANK0/1=AVG/2=MIN/3=MAX>] \n\t" "[-q,--delay ] \n\t" "[-h,--help]\n",