diff --git a/openmp/libomptarget/DeviceRTL/include/Types.h b/openmp/libomptarget/DeviceRTL/include/Types.h --- a/openmp/libomptarget/DeviceRTL/include/Types.h +++ b/openmp/libomptarget/DeviceRTL/include/Types.h @@ -136,7 +136,7 @@ #pragma omp end declare variant #pragma omp begin declare variant match( \ - device = {arch(amdgcn)}, implementation = {extension(match_none)}) + device = {arch(amdgcn)}, implementation = {extension(match_none)}) using LaneMaskTy = uint64_t; #pragma omp end declare variant diff --git a/openmp/libomptarget/DeviceRTL/include/generated_microtask_cases.gen b/openmp/libomptarget/DeviceRTL/include/generated_microtask_cases.gen --- a/openmp/libomptarget/DeviceRTL/include/generated_microtask_cases.gen +++ b/openmp/libomptarget/DeviceRTL/include/generated_microtask_cases.gen @@ -1,405 +1,267 @@ case 0: -((void (*)(int32_t *, int32_t * -))fn)(&global_tid, &bound_tid -); +((void (*)(int32_t *, int32_t *))fn)(&global_tid, &bound_tid); break; case 1: -((void (*)(int32_t *, int32_t * -, void *))fn)(&global_tid, &bound_tid -, args[0]); +((void (*)(int32_t *, int32_t *, void *))fn)(&global_tid, &bound_tid, args[0]); break; case 2: -((void (*)(int32_t *, int32_t * -, void *, void *))fn)(&global_tid, &bound_tid -, args[0], args[1]); +((void (*)(int32_t *, int32_t *, void *, void *))fn)(&global_tid, &bound_tid, + args[0], args[1]); break; case 3: -((void (*)(int32_t *, int32_t * -, void *, void *, void *))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2]); +((void (*)(int32_t *, int32_t *, void *, void *, + void *))fn)(&global_tid, &bound_tid, args[0], args[1], args[2]); break; case 4: -((void (*)(int32_t *, int32_t * -, void *, void *, void *, void * -))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2], args[3] -); +((void (*)(int32_t *, int32_t *, void *, void *, void *, + void *))fn)(&global_tid, &bound_tid, args[0], args[1], args[2], + args[3]); break; case 5: -((void (*)(int32_t *, int32_t * -, void *, void *, void *, void * -, void *))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2], args[3] -, args[4]); +((void (*)(int32_t *, int32_t *, void *, void *, void *, void *, + void *))fn)(&global_tid, &bound_tid, args[0], args[1], args[2], + args[3], args[4]); break; case 6: -((void (*)(int32_t *, int32_t * -, void *, void *, void *, void * -, void *, void *))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2], args[3] -, args[4], args[5]); +((void (*)(int32_t *, int32_t *, void *, void *, void *, void *, void *, + void *))fn)(&global_tid, &bound_tid, args[0], args[1], args[2], + args[3], args[4], args[5]); break; case 7: -((void (*)(int32_t *, int32_t * -, void *, void *, void *, void * -, void *, void *, void *))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2], args[3] -, args[4], args[5], args[6]); +((void (*)(int32_t *, int32_t *, void *, void *, void *, void *, void *, void *, + void *))fn)(&global_tid, &bound_tid, args[0], args[1], args[2], + args[3], args[4], args[5], args[6]); break; case 8: -((void (*)(int32_t *, int32_t * -, void *, void *, void *, void * -, void *, void *, void *, void * -))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2], args[3] -, args[4], args[5], args[6], args[7] -); +((void (*)(int32_t *, int32_t *, void *, void *, void *, void *, void *, void *, + void *, void *))fn)(&global_tid, &bound_tid, args[0], args[1], + args[2], args[3], args[4], args[5], args[6], + args[7]); break; case 9: -((void (*)(int32_t *, int32_t * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2], args[3] -, args[4], args[5], args[6], args[7] -, args[8]); +((void (*)(int32_t *, int32_t *, void *, void *, void *, void *, void *, void *, + void *, void *, void *))fn)(&global_tid, &bound_tid, args[0], + args[1], args[2], args[3], args[4], + args[5], args[6], args[7], args[8]); break; case 10: -((void (*)(int32_t *, int32_t * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2], args[3] -, args[4], args[5], args[6], args[7] -, args[8], args[9]); +((void (*)(int32_t *, int32_t *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *))fn)(&global_tid, &bound_tid, args[0], + args[1], args[2], args[3], + args[4], args[5], args[6], + args[7], args[8], args[9]); break; case 11: -((void (*)(int32_t *, int32_t * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2], args[3] -, args[4], args[5], args[6], args[7] -, args[8], args[9], args[10]); +((void (*)(int32_t *, int32_t *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, + void *))fn)(&global_tid, &bound_tid, args[0], args[1], args[2], + args[3], args[4], args[5], args[6], args[7], args[8], + args[9], args[10]); break; case 12: -((void (*)(int32_t *, int32_t * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2], args[3] -, args[4], args[5], args[6], args[7] -, args[8], args[9], args[10], args[11] -); +((void (*)(int32_t *, int32_t *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, + void *))fn)(&global_tid, &bound_tid, args[0], args[1], args[2], + args[3], args[4], args[5], args[6], args[7], args[8], + args[9], args[10], args[11]); break; case 13: -((void (*)(int32_t *, int32_t * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2], args[3] -, args[4], args[5], args[6], args[7] -, args[8], args[9], args[10], args[11] -, args[12]); +((void (*)(int32_t *, int32_t *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, + void *))fn)(&global_tid, &bound_tid, args[0], args[1], args[2], + args[3], args[4], args[5], args[6], args[7], args[8], + args[9], args[10], args[11], args[12]); break; case 14: -((void (*)(int32_t *, int32_t * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2], args[3] -, args[4], args[5], args[6], args[7] -, args[8], args[9], args[10], args[11] -, args[12], args[13]); +((void (*)(int32_t *, int32_t *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, + void *))fn)(&global_tid, &bound_tid, args[0], args[1], args[2], + args[3], args[4], args[5], args[6], args[7], args[8], + args[9], args[10], args[11], args[12], args[13]); break; case 15: -((void (*)(int32_t *, int32_t * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2], args[3] -, args[4], args[5], args[6], args[7] -, args[8], args[9], args[10], args[11] -, args[12], args[13], args[14]); +((void (*)(int32_t *, int32_t *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *))fn)(&global_tid, &bound_tid, args[0], args[1], args[2], + args[3], args[4], args[5], args[6], args[7], args[8], + args[9], args[10], args[11], args[12], args[13], + args[14]); break; case 16: -((void (*)(int32_t *, int32_t * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2], args[3] -, args[4], args[5], args[6], args[7] -, args[8], args[9], args[10], args[11] -, args[12], args[13], args[14], args[15] -); +((void (*)(int32_t *, int32_t *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *, void *))fn)(&global_tid, &bound_tid, args[0], args[1], + args[2], args[3], args[4], args[5], args[6], + args[7], args[8], args[9], args[10], args[11], + args[12], args[13], args[14], args[15]); break; case 17: -((void (*)(int32_t *, int32_t * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2], args[3] -, args[4], args[5], args[6], args[7] -, args[8], args[9], args[10], args[11] -, args[12], args[13], args[14], args[15] -, args[16]); +((void (*)(int32_t *, int32_t *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *, void *, void *))fn)(&global_tid, &bound_tid, args[0], + args[1], args[2], args[3], args[4], + args[5], args[6], args[7], args[8], + args[9], args[10], args[11], args[12], + args[13], args[14], args[15], args[16]); break; case 18: -((void (*)(int32_t *, int32_t * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2], args[3] -, args[4], args[5], args[6], args[7] -, args[8], args[9], args[10], args[11] -, args[12], args[13], args[14], args[15] -, args[16], args[17]); +((void (*)(int32_t *, int32_t *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, + void *))fn)(&global_tid, &bound_tid, args[0], args[1], args[2], + args[3], args[4], args[5], args[6], args[7], args[8], + args[9], args[10], args[11], args[12], args[13], + args[14], args[15], args[16], args[17]); break; case 19: -((void (*)(int32_t *, int32_t * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2], args[3] -, args[4], args[5], args[6], args[7] -, args[8], args[9], args[10], args[11] -, args[12], args[13], args[14], args[15] -, args[16], args[17], args[18]); +((void (*)(int32_t *, int32_t *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, + void *))fn)(&global_tid, &bound_tid, args[0], args[1], args[2], + args[3], args[4], args[5], args[6], args[7], args[8], + args[9], args[10], args[11], args[12], args[13], + args[14], args[15], args[16], args[17], args[18]); break; case 20: -((void (*)(int32_t *, int32_t * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2], args[3] -, args[4], args[5], args[6], args[7] -, args[8], args[9], args[10], args[11] -, args[12], args[13], args[14], args[15] -, args[16], args[17], args[18], args[19] -); +((void (*)(int32_t *, int32_t *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *))fn)( + &global_tid, &bound_tid, args[0], args[1], args[2], args[3], args[4], + args[5], args[6], args[7], args[8], args[9], args[10], args[11], args[12], + args[13], args[14], args[15], args[16], args[17], args[18], args[19]); break; case 21: -((void (*)(int32_t *, int32_t * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2], args[3] -, args[4], args[5], args[6], args[7] -, args[8], args[9], args[10], args[11] -, args[12], args[13], args[14], args[15] -, args[16], args[17], args[18], args[19] -, args[20]); +((void (*)(int32_t *, int32_t *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, + void *))fn)(&global_tid, &bound_tid, args[0], args[1], args[2], + args[3], args[4], args[5], args[6], args[7], args[8], + args[9], args[10], args[11], args[12], args[13], + args[14], args[15], args[16], args[17], args[18], + args[19], args[20]); break; case 22: -((void (*)(int32_t *, int32_t * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2], args[3] -, args[4], args[5], args[6], args[7] -, args[8], args[9], args[10], args[11] -, args[12], args[13], args[14], args[15] -, args[16], args[17], args[18], args[19] -, args[20], args[21]); +((void (*)(int32_t *, int32_t *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, + void *))fn)(&global_tid, &bound_tid, args[0], args[1], args[2], + args[3], args[4], args[5], args[6], args[7], args[8], + args[9], args[10], args[11], args[12], args[13], + args[14], args[15], args[16], args[17], args[18], + args[19], args[20], args[21]); break; case 23: -((void (*)(int32_t *, int32_t * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2], args[3] -, args[4], args[5], args[6], args[7] -, args[8], args[9], args[10], args[11] -, args[12], args[13], args[14], args[15] -, args[16], args[17], args[18], args[19] -, args[20], args[21], args[22]); +((void (*)(int32_t *, int32_t *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *))fn)(&global_tid, &bound_tid, args[0], args[1], args[2], + args[3], args[4], args[5], args[6], args[7], args[8], + args[9], args[10], args[11], args[12], args[13], + args[14], args[15], args[16], args[17], args[18], + args[19], args[20], args[21], args[22]); break; case 24: -((void (*)(int32_t *, int32_t * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2], args[3] -, args[4], args[5], args[6], args[7] -, args[8], args[9], args[10], args[11] -, args[12], args[13], args[14], args[15] -, args[16], args[17], args[18], args[19] -, args[20], args[21], args[22], args[23] -); +((void (*)(int32_t *, int32_t *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *, void *))fn)(&global_tid, &bound_tid, args[0], args[1], + args[2], args[3], args[4], args[5], args[6], + args[7], args[8], args[9], args[10], args[11], + args[12], args[13], args[14], args[15], args[16], + args[17], args[18], args[19], args[20], args[21], + args[22], args[23]); break; case 25: -((void (*)(int32_t *, int32_t * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2], args[3] -, args[4], args[5], args[6], args[7] -, args[8], args[9], args[10], args[11] -, args[12], args[13], args[14], args[15] -, args[16], args[17], args[18], args[19] -, args[20], args[21], args[22], args[23] -, args[24]); +((void (*)(int32_t *, int32_t *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *, void *, void *))fn)(&global_tid, &bound_tid, args[0], + args[1], args[2], args[3], args[4], + args[5], args[6], args[7], args[8], + args[9], args[10], args[11], args[12], + args[13], args[14], args[15], args[16], + args[17], args[18], args[19], args[20], + args[21], args[22], args[23], args[24]); break; case 26: -((void (*)(int32_t *, int32_t * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2], args[3] -, args[4], args[5], args[6], args[7] -, args[8], args[9], args[10], args[11] -, args[12], args[13], args[14], args[15] -, args[16], args[17], args[18], args[19] -, args[20], args[21], args[22], args[23] -, args[24], args[25]); +((void (*)(int32_t *, int32_t *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *))fn)( + &global_tid, &bound_tid, args[0], args[1], args[2], args[3], args[4], + args[5], args[6], args[7], args[8], args[9], args[10], args[11], args[12], + args[13], args[14], args[15], args[16], args[17], args[18], args[19], + args[20], args[21], args[22], args[23], args[24], args[25]); break; case 27: -((void (*)(int32_t *, int32_t * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2], args[3] -, args[4], args[5], args[6], args[7] -, args[8], args[9], args[10], args[11] -, args[12], args[13], args[14], args[15] -, args[16], args[17], args[18], args[19] -, args[20], args[21], args[22], args[23] -, args[24], args[25], args[26]); +((void (*)(int32_t *, int32_t *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *))fn)( + &global_tid, &bound_tid, args[0], args[1], args[2], args[3], args[4], + args[5], args[6], args[7], args[8], args[9], args[10], args[11], args[12], + args[13], args[14], args[15], args[16], args[17], args[18], args[19], + args[20], args[21], args[22], args[23], args[24], args[25], args[26]); break; case 28: -((void (*)(int32_t *, int32_t * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2], args[3] -, args[4], args[5], args[6], args[7] -, args[8], args[9], args[10], args[11] -, args[12], args[13], args[14], args[15] -, args[16], args[17], args[18], args[19] -, args[20], args[21], args[22], args[23] -, args[24], args[25], args[26], args[27] -); +((void (*)(int32_t *, int32_t *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, + void *))fn)(&global_tid, &bound_tid, args[0], args[1], args[2], + args[3], args[4], args[5], args[6], args[7], args[8], + args[9], args[10], args[11], args[12], args[13], + args[14], args[15], args[16], args[17], args[18], + args[19], args[20], args[21], args[22], args[23], + args[24], args[25], args[26], args[27]); break; case 29: -((void (*)(int32_t *, int32_t * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2], args[3] -, args[4], args[5], args[6], args[7] -, args[8], args[9], args[10], args[11] -, args[12], args[13], args[14], args[15] -, args[16], args[17], args[18], args[19] -, args[20], args[21], args[22], args[23] -, args[24], args[25], args[26], args[27] -, args[28]); +((void (*)(int32_t *, int32_t *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, + void *))fn)(&global_tid, &bound_tid, args[0], args[1], args[2], + args[3], args[4], args[5], args[6], args[7], args[8], + args[9], args[10], args[11], args[12], args[13], + args[14], args[15], args[16], args[17], args[18], + args[19], args[20], args[21], args[22], args[23], + args[24], args[25], args[26], args[27], args[28]); break; case 30: -((void (*)(int32_t *, int32_t * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2], args[3] -, args[4], args[5], args[6], args[7] -, args[8], args[9], args[10], args[11] -, args[12], args[13], args[14], args[15] -, args[16], args[17], args[18], args[19] -, args[20], args[21], args[22], args[23] -, args[24], args[25], args[26], args[27] -, args[28], args[29]); +((void (*)(int32_t *, int32_t *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *))fn)( + &global_tid, &bound_tid, args[0], args[1], args[2], args[3], args[4], + args[5], args[6], args[7], args[8], args[9], args[10], args[11], args[12], + args[13], args[14], args[15], args[16], args[17], args[18], args[19], + args[20], args[21], args[22], args[23], args[24], args[25], args[26], + args[27], args[28], args[29]); break; case 31: -((void (*)(int32_t *, int32_t * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2], args[3] -, args[4], args[5], args[6], args[7] -, args[8], args[9], args[10], args[11] -, args[12], args[13], args[14], args[15] -, args[16], args[17], args[18], args[19] -, args[20], args[21], args[22], args[23] -, args[24], args[25], args[26], args[27] -, args[28], args[29], args[30]); +((void (*)(int32_t *, int32_t *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *))fn)(&global_tid, &bound_tid, args[0], args[1], args[2], + args[3], args[4], args[5], args[6], args[7], args[8], + args[9], args[10], args[11], args[12], args[13], + args[14], args[15], args[16], args[17], args[18], + args[19], args[20], args[21], args[22], args[23], + args[24], args[25], args[26], args[27], args[28], + args[29], args[30]); break; case 32: -((void (*)(int32_t *, int32_t * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -, void *, void *, void *, void * -))fn)(&global_tid, &bound_tid -, args[0], args[1], args[2], args[3] -, args[4], args[5], args[6], args[7] -, args[8], args[9], args[10], args[11] -, args[12], args[13], args[14], args[15] -, args[16], args[17], args[18], args[19] -, args[20], args[21], args[22], args[23] -, args[24], args[25], args[26], args[27] -, args[28], args[29], args[30], args[31] -); +((void (*)(int32_t *, int32_t *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *, void *, void *, void *, void *, void *, void *, void *, + void *, void *))fn)(&global_tid, &bound_tid, args[0], args[1], + args[2], args[3], args[4], args[5], args[6], + args[7], args[8], args[9], args[10], args[11], + args[12], args[13], args[14], args[15], args[16], + args[17], args[18], args[19], args[20], args[21], + args[22], args[23], args[24], args[25], args[26], + args[27], args[28], args[29], args[30], + args[31]); break; diff --git a/openmp/libomptarget/DeviceRTL/src/LibC.cpp b/openmp/libomptarget/DeviceRTL/src/LibC.cpp --- a/openmp/libomptarget/DeviceRTL/src/LibC.cpp +++ b/openmp/libomptarget/DeviceRTL/src/LibC.cpp @@ -15,7 +15,8 @@ } #pragma omp begin declare variant match( \ - device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)}) + device = {arch(nvptx, nvptx64)}, \ + implementation = {extension(match_any)}) extern "C" int32_t vprintf(const char *, void *); namespace impl { int32_t omp_vprintf(const char *Format, void *Arguments, uint32_t) { diff --git a/openmp/libomptarget/DeviceRTL/src/Mapping.cpp b/openmp/libomptarget/DeviceRTL/src/Mapping.cpp --- a/openmp/libomptarget/DeviceRTL/src/Mapping.cpp +++ b/openmp/libomptarget/DeviceRTL/src/Mapping.cpp @@ -98,7 +98,8 @@ /// ///{ #pragma omp begin declare variant match( \ - device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)}) + device = {arch(nvptx, nvptx64)}, \ + implementation = {extension(match_any)}) uint32_t getNumHardwareThreadsInBlock() { return __nvvm_read_ptx_sreg_ntid_x(); diff --git a/openmp/libomptarget/DeviceRTL/src/Misc.cpp b/openmp/libomptarget/DeviceRTL/src/Misc.cpp --- a/openmp/libomptarget/DeviceRTL/src/Misc.cpp +++ b/openmp/libomptarget/DeviceRTL/src/Misc.cpp @@ -43,7 +43,8 @@ /// ///{ #pragma omp begin declare variant match( \ - device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)}) + device = {arch(nvptx, nvptx64)}, \ + implementation = {extension(match_any)}) double getWTick() { // Timer precision is 1ns diff --git a/openmp/libomptarget/DeviceRTL/src/Synchronization.cpp b/openmp/libomptarget/DeviceRTL/src/Synchronization.cpp --- a/openmp/libomptarget/DeviceRTL/src/Synchronization.cpp +++ b/openmp/libomptarget/DeviceRTL/src/Synchronization.cpp @@ -291,7 +291,7 @@ if (mapping::getThreadIdInWarp() == LowestActiveThread) { fenceKernel(atomic::release); while (!atomicCAS((uint32_t *)Lock, UNSET, SET, atomic::relaxed, - atomic::relaxed)) { + atomic::relaxed)) { __builtin_amdgcn_s_sleep(32); } fenceKernel(atomic::aquire); @@ -305,7 +305,8 @@ /// ///{ #pragma omp begin declare variant match( \ - device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)}) + device = {arch(nvptx, nvptx64)}, \ + implementation = {extension(match_any)}) uint32_t atomicInc(uint32_t *Address, uint32_t Val, atomic::OrderingTy Ordering) { @@ -483,13 +484,9 @@ return impl::atomicInc(Addr, V, Ordering); } -void unsetCriticalLock(omp_lock_t *Lock) { - impl::unsetLock(Lock); -} +void unsetCriticalLock(omp_lock_t *Lock) { impl::unsetLock(Lock); } -void setCriticalLock(omp_lock_t *Lock) { - impl::setLock(Lock); -} +void setCriticalLock(omp_lock_t *Lock) { impl::setLock(Lock); } extern "C" { void __kmpc_ordered(IdentTy *Loc, int32_t TId) { FunctionTracingRAII(); } diff --git a/openmp/libomptarget/DeviceRTL/src/Utils.cpp b/openmp/libomptarget/DeviceRTL/src/Utils.cpp --- a/openmp/libomptarget/DeviceRTL/src/Utils.cpp +++ b/openmp/libomptarget/DeviceRTL/src/Utils.cpp @@ -59,7 +59,8 @@ /// ///{ #pragma omp begin declare variant match( \ - device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)}) + device = {arch(nvptx, nvptx64)}, \ + implementation = {extension(match_any)}) void Unpack(uint64_t Val, uint32_t *LowBits, uint32_t *HighBits) { uint32_t LowBitsLocal, HighBitsLocal; @@ -103,8 +104,9 @@ return __builtin_amdgcn_ds_bpermute(Index << 2, Var); } -bool isSharedMemPtr(const void * Ptr) { - return __builtin_amdgcn_is_shared((const __attribute__((address_space(0))) void *)Ptr); +bool isSharedMemPtr(const void *Ptr) { + return __builtin_amdgcn_is_shared( + (const __attribute__((address_space(0))) void *)Ptr); } #pragma omp end declare variant ///} @@ -113,7 +115,8 @@ /// ///{ #pragma omp begin declare variant match( \ - device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)}) + device = {arch(nvptx, nvptx64)}, \ + implementation = {extension(match_any)}) int32_t shuffle(uint64_t Mask, int32_t Var, int32_t SrcLane) { return __nvvm_shfl_sync_idx_i32(Mask, Var, SrcLane, 0x1f); diff --git a/openmp/libomptarget/DeviceRTL/src/Workshare.cpp b/openmp/libomptarget/DeviceRTL/src/Workshare.cpp --- a/openmp/libomptarget/DeviceRTL/src/Workshare.cpp +++ b/openmp/libomptarget/DeviceRTL/src/Workshare.cpp @@ -114,9 +114,9 @@ //////////////////////////////////////////////////////////////////////////////// // Support for Static Init - static void for_static_init(int32_t, int32_t schedtype, - int32_t *plastiter, T *plower, T *pupper, - ST *pstride, ST chunk, bool IsSPMDExecutionMode) { + static void for_static_init(int32_t, int32_t schedtype, int32_t *plastiter, + T *plower, T *pupper, ST *pstride, ST chunk, + bool IsSPMDExecutionMode) { int32_t gtid = omp_get_thread_num(); int numberOfActiveOMPThreads = omp_get_num_threads(); diff --git a/openmp/libomptarget/DeviceRTL/src/exports b/openmp/libomptarget/DeviceRTL/src/exports --- a/openmp/libomptarget/DeviceRTL/src/exports +++ b/openmp/libomptarget/DeviceRTL/src/exports @@ -1,12 +1,7 @@ -omp_* -*llvm_* -__kmpc_* +omp_ **llvm_ *__kmpc_ * -_ZN4ompx* + _ZN4ompx * -__keep_alive -IsSPMDMode + __keep_alive IsSPMDMode -memcmp -printf -__assert_fail + memcmp printf __assert_fail