[OpenMP][NFC] clang-format the OpenMP device runtime

These files aren't fully formatted. I'm guessing this was a holdover
from when `clang-format` was totally broken for OpenMP offloading.
Format the files to be more consistent.

Reviewed By: tianshilei1992

Differential Revision: https://reviews.llvm.org/D151226
This commit is contained in:
Joseph Huber 2023-05-23 11:09:16 -05:00
parent e826762a08
commit 47800a12dc
8 changed files with 223 additions and 358 deletions

View File

@ -136,7 +136,7 @@ using LaneMaskTy = uint64_t;
#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

View File

@ -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;

View File

@ -15,7 +15,8 @@ int32_t omp_vprintf(const char *Format, void *Arguments, uint32_t);
}
#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) {

View File

@ -98,7 +98,8 @@ uint32_t getNumberOfWarpsInBlock() {
///
///{
#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();

View File

@ -43,7 +43,8 @@ double getWTime() {
///
///{
#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

View File

@ -291,7 +291,7 @@ void setCriticalLock(omp_lock_t *Lock) {
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 @@ void setCriticalLock(omp_lock_t *Lock) {
///
///{
#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 @@ uint32_t atomic::inc(uint32_t *Addr, uint32_t V, atomic::OrderingTy Ordering) {
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(); }

View File

@ -59,7 +59,8 @@ uint64_t Pack(uint32_t LowBits, uint32_t HighBits) {
///
///{
#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 @@ int32_t shuffleDown(uint64_t Mask, int32_t Var, uint32_t LaneDelta,
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 @@ bool isSharedMemPtr(const void * Ptr) {
///
///{
#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);

View File

@ -114,9 +114,9 @@ template <typename T, typename ST> struct omptarget_nvptx_LoopSupport {
////////////////////////////////////////////////////////////////////////////////
// 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();