AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Cache flush

Browse Source
This commit is contained in:
saurabhAMD
2024-05-07 11:09:32 -05:00
parent 16dfeaf89b
commit 3c0728e8eb
2 changed files with 53 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/common.cu
+52 -1
View File
@@ -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 <num graph launches>] \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 <delay between out-of-place and in-place in microseconds>] \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;
src/common.h
+1
View File
@@ -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;