This is an archive of the discontinued LLVM Phabricator instance.

Differential D151226

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

Authored by jhuber6 on May 23 2023, 9:10 AM.

Details

Reviewers
JonChesterfield
tianshilei1992
jdoerfert
Commits
rG47800a12dcb4: [OpenMP][NFC] clang-format the OpenMP device runtime
Summary

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.

Diff Detail

Event Timeline

jhuber6 created this revision.May 23 2023, 9:10 AM
Herald added a project: Restricted Project. · View Herald TranscriptMay 23 2023, 9:10 AM
Herald added subscribers: sunshaoce, guansong, yaxunl. · View Herald Transcript
jhuber6 requested review of this revision.May 23 2023, 9:10 AM
Herald added a project: Restricted Project. · View Herald TranscriptMay 23 2023, 9:10 AM
Herald added subscribers: openmp-commits, jplehr, sstefan1. · View Herald Transcript
jhuber6 updated this revision to Diff 524757.May 23 2023, 9:11 AM

Exports file got formatted accidentally.

Harbormaster completed remote builds in B233902: Diff 524757.May 23 2023, 9:14 AM
tianshilei1992 accepted this revision.May 23 2023, 9:18 AM

LG

This revision is now accepted and ready to land.May 23 2023, 9:18 AM
Closed by commit rG47800a12dcb4: [OpenMP][NFC] clang-format the OpenMP device runtime (authored by jhuber6). · Explain WhyMay 23 2023, 9:19 AM
This revision was automatically updated to reflect the committed changes.
jhuber6 added a commit: rG47800a12dcb4: [OpenMP][NFC] clang-format the OpenMP device runtime.
protze.joachim added a subscriber: protze.joachim.May 24 2023, 7:53 AM

If anyone cares for usable git blame functionality across reformatting commits, we should add the openmp clang-format commits to the top-level .git-blame-ignore-revs file. Only commits with no further changes than formatting should be added to the file.
The file contains the command to configure the local git repo. GitHub respects this file for online blame view.

jhuber6 added a comment.May 24 2023, 7:57 AM
In D151226#4368522, @protze.joachim wrote:

If anyone cares for usable git blame functionality across reformatting commits, we should add the openmp clang-format commits to the top-level .git-blame-ignore-revs file. Only commits with no further changes than formatting should be added to the file.
The file contains the command to configure the local git repo. GitHub respects this file for online blame view.

Didn't know about that, I'll keep it in mind. Thanks.

protze.joachim added a comment.May 24 2023, 9:00 AM

Since you need to know the commit hash, adding the formatting commit to the file is always a separate commit.

Revision Contents

PathSize
openmp/
libomptarget/
DeviceRTL/
include/
2 lines
540 lines
src/
3 lines
3 lines
3 lines
13 lines
11 lines
6 lines

Diff 524761

openmp/libomptarget/DeviceRTL/include/Types.h

Show First 20 Lines • Show All 130 Lines • ▼ Show 20 Linesstruct TaskDescriptorTy {
TaskFnTy TaskFn; TaskFnTy TaskFn;
}; };
#pragma omp begin declare variant match(device = {arch(amdgcn)}) #pragma omp begin declare variant match(device = {arch(amdgcn)})
using LaneMaskTy = uint64_t; using LaneMaskTy = uint64_t;
#pragma omp end declare variant #pragma omp end declare variant
#pragma omp begin declare variant match( \ #pragma omp begin declare variant match( \
device = {arch(amdgcn)}, implementation = {extension(match_none)}) device = {arch(amdgcn)}, implementation = {extension(match_none)})
using LaneMaskTy = uint64_t; using LaneMaskTy = uint64_t;
#pragma omp end declare variant #pragma omp end declare variant
namespace lanes { namespace lanes {
enum : LaneMaskTy { All = ~(LaneMaskTy)0 }; enum : LaneMaskTy { All = ~(LaneMaskTy)0 };
} // namespace lanes } // namespace lanes
/// The ident structure that describes a source location. The struct is /// The ident structure that describes a source location. The struct is
▲ Show 20 Lines • Show All 66 Lines • Show Last 20 Lines

openmp/libomptarget/DeviceRTL/include/generated_microtask_cases.gen

case 0: case 0:
((void (*)(int32_t *, int32_t * ((void (*)(int32_t *, int32_t *))fn)(&global_tid, &bound_tid);
))fn)(&global_tid, &bound_tid
);
break; break;
case 1: case 1:
((void (*)(int32_t *, int32_t * ((void (*)(int32_t *, int32_t *, void *))fn)(&global_tid, &bound_tid, args[0]);
, void *))fn)(&global_tid, &bound_tid
, args[0]);
break; break;
case 2: case 2:
((void (*)(int32_t *, int32_t * ((void (*)(int32_t *, int32_t *, void *, void *))fn)(&global_tid, &bound_tid,
, void *, void *))fn)(&global_tid, &bound_tid args[0], args[1]);
, args[0], args[1]);
break; break;
case 3: case 3:
((void (*)(int32_t *, int32_t * ((void (*)(int32_t *, int32_t *, void *, void *,
, void *, void *, void *))fn)(&global_tid, &bound_tid void *))fn)(&global_tid, &bound_tid, args[0], args[1], args[2]);
, args[0], args[1], args[2]);
break; break;
case 4: case 4:
((void (*)(int32_t *, int32_t * ((void (*)(int32_t *, int32_t *, void *, void *, void *,
, void *, void *, void *, void * void *))fn)(&global_tid, &bound_tid, args[0], args[1], args[2],
))fn)(&global_tid, &bound_tid args[3]);
, args[0], args[1], args[2], args[3]
);
break; break;
case 5: case 5:
((void (*)(int32_t *, int32_t * ((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],
, void *))fn)(&global_tid, &bound_tid args[3], args[4]);
, args[0], args[1], args[2], args[3]
, args[4]);
break; break;
case 6: case 6:
((void (*)(int32_t *, int32_t * ((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],
, void *, void *))fn)(&global_tid, &bound_tid args[3], args[4], args[5]);
, args[0], args[1], args[2], args[3]
, args[4], args[5]);
break; break;
case 7: case 7:
((void (*)(int32_t *, int32_t * ((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],
, void *, void *, void *))fn)(&global_tid, &bound_tid args[3], args[4], args[5], args[6]);
, args[0], args[1], args[2], args[3]
, args[4], args[5], args[6]);
break; break;
case 8: case 8:
((void (*)(int32_t *, int32_t * ((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],
, void *, void *, void *, void * args[2], args[3], args[4], args[5], args[6],
))fn)(&global_tid, &bound_tid args[7]);
, args[0], args[1], args[2], args[3]
, args[4], args[5], args[6], args[7]
);
break; break;
case 9: case 9:
((void (*)(int32_t *, int32_t * ((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],
, void *, void *, void *, void * args[1], args[2], args[3], args[4],
, void *))fn)(&global_tid, &bound_tid args[5], args[6], args[7], args[8]);
, args[0], args[1], args[2], args[3]
, args[4], args[5], args[6], args[7]
, args[8]);
break; break;
case 10: case 10:
((void (*)(int32_t *, int32_t * ((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],
, void *, void *, void *, void * args[1], args[2], args[3],
, void *, void *))fn)(&global_tid, &bound_tid args[4], args[5], args[6],
, args[0], args[1], args[2], args[3] args[7], args[8], args[9]);
, args[4], args[5], args[6], args[7]
, args[8], args[9]);
break; break;
case 11: case 11:
((void (*)(int32_t *, int32_t * ((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],
, void *, void *, void *))fn)(&global_tid, &bound_tid args[3], args[4], args[5], args[6], args[7], args[8],
, args[0], args[1], args[2], args[3] args[9], args[10]);
, args[4], args[5], args[6], args[7]
, args[8], args[9], args[10]);
break; break;
case 12: case 12:
((void (*)(int32_t *, int32_t * ((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],
, void *, void *, void *, void * args[3], args[4], args[5], args[6], args[7], args[8],
))fn)(&global_tid, &bound_tid args[9], args[10], args[11]);
, 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; break;
case 13: case 13:
((void (*)(int32_t *, int32_t * ((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],
, void *, void *, void *, void * args[3], args[4], args[5], args[6], args[7], args[8],
, void *))fn)(&global_tid, &bound_tid args[9], args[10], args[11], args[12]);
, 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; break;
case 14: case 14:
((void (*)(int32_t *, int32_t * ((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],
, void *, void *, void *, void * args[3], args[4], args[5], args[6], args[7], args[8],
, void *, void *))fn)(&global_tid, &bound_tid args[9], args[10], args[11], args[12], args[13]);
, 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; break;
case 15: case 15:
((void (*)(int32_t *, int32_t * ((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],
, void *, void *, void *, void * args[3], args[4], args[5], args[6], args[7], args[8],
, void *, void *, void *))fn)(&global_tid, &bound_tid args[9], args[10], args[11], args[12], args[13],
, args[0], args[1], args[2], args[3] args[14]);
, args[4], args[5], args[6], args[7]
, args[8], args[9], args[10], args[11]
, args[12], args[13], args[14]);
break; break;
case 16: case 16:
((void (*)(int32_t *, int32_t * ((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],
, void *, void *, void *, void * args[2], args[3], args[4], args[5], args[6],
, void *, void *, void *, void * args[7], args[8], args[9], args[10], args[11],
))fn)(&global_tid, &bound_tid args[12], args[13], args[14], args[15]);
, 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; break;
case 17: case 17:
((void (*)(int32_t *, int32_t * ((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],
, void *, void *, void *, void * args[1], args[2], args[3], args[4],
, void *, void *, void *, void * args[5], args[6], args[7], args[8],
, void *))fn)(&global_tid, &bound_tid args[9], args[10], args[11], args[12],
, args[0], args[1], args[2], args[3] args[13], args[14], args[15], args[16]);
, 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; break;
case 18: case 18:
((void (*)(int32_t *, int32_t * ((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],
, void *, void *, void *, void * args[3], args[4], args[5], args[6], args[7], args[8],
, void *, void *))fn)(&global_tid, &bound_tid args[9], args[10], args[11], args[12], args[13],
, args[0], args[1], args[2], args[3] args[14], args[15], args[16], args[17]);
, 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; break;
case 19: case 19:
((void (*)(int32_t *, int32_t * ((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],
, void *, void *, void *, void * args[3], args[4], args[5], args[6], args[7], args[8],
, void *, void *, void *))fn)(&global_tid, &bound_tid args[9], args[10], args[11], args[12], args[13],
, args[0], args[1], args[2], args[3] args[14], args[15], args[16], args[17], args[18]);
, 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; break;
case 20: case 20:
((void (*)(int32_t *, int32_t * ((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)(
, void *, void *, void *, void * &global_tid, &bound_tid, args[0], args[1], args[2], args[3], args[4],
, void *, void *, void *, void * args[5], args[6], args[7], args[8], args[9], args[10], args[11], args[12],
, void *, void *, void *, void * args[13], args[14], args[15], args[16], args[17], args[18], args[19]);
))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; break;
case 21: case 21:
((void (*)(int32_t *, int32_t * ((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 * void *))fn)(&global_tid, &bound_tid, args[0], args[1], args[2],
, void *, void *, void *, void * args[3], args[4], args[5], args[6], args[7], args[8],
, void *, void *, void *, void * args[9], args[10], args[11], args[12], args[13],
, void *))fn)(&global_tid, &bound_tid args[14], args[15], args[16], args[17], args[18],
, args[0], args[1], args[2], args[3] args[19], args[20]);
, 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; break;
case 22: case 22:
((void (*)(int32_t *, int32_t * ((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 *, void * void *))fn)(&global_tid, &bound_tid, args[0], args[1], args[2],
, void *, void *, void *, void * args[3], args[4], args[5], args[6], args[7], args[8],
, void *, void *, void *, void * args[9], args[10], args[11], args[12], args[13],
, void *, void *))fn)(&global_tid, &bound_tid args[14], args[15], args[16], args[17], args[18],
, args[0], args[1], args[2], args[3] args[19], args[20], args[21]);
, 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; break;
case 23: case 23:
((void (*)(int32_t *, int32_t * ((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 *, void *, void * void *))fn)(&global_tid, &bound_tid, args[0], args[1], args[2],
, void *, void *, void *, void * args[3], args[4], args[5], args[6], args[7], args[8],
, void *, void *, void *, void * args[9], args[10], args[11], args[12], args[13],
, void *, void *, void *))fn)(&global_tid, &bound_tid args[14], args[15], args[16], args[17], args[18],
, args[0], args[1], args[2], args[3] args[19], args[20], args[21], args[22]);
, 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; break;
case 24: case 24:
((void (*)(int32_t *, int32_t * ((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 *, void *, void * void *, void *))fn)(&global_tid, &bound_tid, args[0], args[1],
, void *, void *, void *, void * args[2], args[3], args[4], args[5], args[6],
, void *, void *, void *, void * args[7], args[8], args[9], args[10], args[11],
, void *, void *, void *, void * args[12], args[13], args[14], args[15], args[16],
))fn)(&global_tid, &bound_tid args[17], args[18], args[19], args[20], args[21],
, args[0], args[1], args[2], args[3] args[22], args[23]);
, 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; break;
case 25: case 25:
((void (*)(int32_t *, int32_t * ((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 *, void *, void * void *, void *, void *))fn)(&global_tid, &bound_tid, args[0],
, void *, void *, void *, void * args[1], args[2], args[3], args[4],
, void *, void *, void *, void * args[5], args[6], args[7], args[8],
, void *, void *, void *, void * args[9], args[10], args[11], args[12],
, void *))fn)(&global_tid, &bound_tid args[13], args[14], args[15], args[16],
, args[0], args[1], args[2], args[3] args[17], args[18], args[19], args[20],
, args[4], args[5], args[6], args[7] args[21], args[22], args[23], args[24]);
, 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; break;
case 26: case 26:
((void (*)(int32_t *, int32_t * ((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 *, void *, void * void *, void *, void *, void *))fn)(
, void *, void *, void *, void * &global_tid, &bound_tid, args[0], args[1], args[2], args[3], args[4],
, void *, void *, void *, void * args[5], args[6], args[7], args[8], args[9], args[10], args[11], args[12],
, void *, void *, void *, void * args[13], args[14], args[15], args[16], args[17], args[18], args[19],
, void *, void *))fn)(&global_tid, &bound_tid args[20], args[21], args[22], args[23], args[24], args[25]);
, 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; break;
case 27: case 27:
((void (*)(int32_t *, int32_t * ((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 *, void *, void * void *, void *, void *, void *, void *))fn)(
, void *, void *, void *, void * &global_tid, &bound_tid, args[0], args[1], args[2], args[3], args[4],
, void *, void *, void *, void * args[5], args[6], args[7], args[8], args[9], args[10], args[11], args[12],
, void *, void *, void *, void * args[13], args[14], args[15], args[16], args[17], args[18], args[19],
, void *, void *, void *))fn)(&global_tid, &bound_tid args[20], args[21], args[22], args[23], args[24], args[25], args[26]);
, 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; break;
case 28: case 28:
((void (*)(int32_t *, int32_t * ((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 *, void *, void * void *, void *, void *, void *, void *,
, void *, void *, void *, void * void *))fn)(&global_tid, &bound_tid, args[0], args[1], args[2],
, void *, void *, void *, void * args[3], args[4], args[5], args[6], args[7], args[8],
, void *, void *, void *, void * args[9], args[10], args[11], args[12], args[13],
, void *, void *, void *, void * args[14], args[15], args[16], args[17], args[18],
))fn)(&global_tid, &bound_tid args[19], args[20], args[21], args[22], args[23],
, args[0], args[1], args[2], args[3] args[24], args[25], args[26], args[27]);
, 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; break;
case 29: case 29:
((void (*)(int32_t *, int32_t * ((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 *, void *, void * void *, void *, void *, void *, void *, void *,
, void *, void *, void *, void * void *))fn)(&global_tid, &bound_tid, args[0], args[1], args[2],
, void *, void *, void *, void * args[3], args[4], args[5], args[6], args[7], args[8],
, void *, void *, void *, void * args[9], args[10], args[11], args[12], args[13],
, void *, void *, void *, void * args[14], args[15], args[16], args[17], args[18],
, void *))fn)(&global_tid, &bound_tid args[19], args[20], args[21], args[22], args[23],
, args[0], args[1], args[2], args[3] args[24], args[25], args[26], args[27], args[28]);
, 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; break;
case 30: case 30:
((void (*)(int32_t *, int32_t * ((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 *, void *, void * void *, void *, void *, void *, void *, void *, void *, void *))fn)(
, void *, void *, void *, void * &global_tid, &bound_tid, args[0], args[1], args[2], args[3], args[4],
, void *, void *, void *, void * args[5], args[6], args[7], args[8], args[9], args[10], args[11], args[12],
, void *, void *, void *, void * args[13], args[14], args[15], args[16], args[17], args[18], args[19],
, void *, void *, void *, void * args[20], args[21], args[22], args[23], args[24], args[25], args[26],
, void *, void *))fn)(&global_tid, &bound_tid args[27], args[28], args[29]);
, 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; break;
case 31: case 31:
((void (*)(int32_t *, int32_t * ((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 *, 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],
, void *, void *, void *, void * args[3], args[4], args[5], args[6], args[7], args[8],
, void *, void *, void *, void * args[9], args[10], args[11], args[12], args[13],
, void *, void *, void *, void * args[14], args[15], args[16], args[17], args[18],
, void *, void *, void *))fn)(&global_tid, &bound_tid args[19], args[20], args[21], args[22], args[23],
, args[0], args[1], args[2], args[3] args[24], args[25], args[26], args[27], args[28],
, args[4], args[5], args[6], args[7] args[29], args[30]);
, 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; break;
case 32: case 32:
((void (*)(int32_t *, int32_t * ((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 *, 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],
, void *, void *, void *, void * args[2], args[3], args[4], args[5], args[6],
, void *, void *, void *, void * args[7], args[8], args[9], args[10], args[11],
, void *, void *, void *, void * args[12], args[13], args[14], args[15], args[16],
, void *, void *, void *, void * args[17], args[18], args[19], args[20], args[21],
))fn)(&global_tid, &bound_tid args[22], args[23], args[24], args[25], args[26],
, args[0], args[1], args[2], args[3] args[27], args[28], args[29], args[30],
, args[4], args[5], args[6], args[7] args[31]);
, 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; break;

openmp/libomptarget/DeviceRTL/src/LibC.cpp

Show All 9 Lines
#pragma omp begin declare target device_type(nohost) #pragma omp begin declare target device_type(nohost)
namespace impl { namespace impl {
int32_t omp_vprintf(const char *Format, void *Arguments, uint32_t); int32_t omp_vprintf(const char *Format, void *Arguments, uint32_t);
} }
#pragma omp begin declare variant match( \ #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 *); extern "C" int32_t vprintf(const char *, void *);
namespace impl { namespace impl {
int32_t omp_vprintf(const char *Format, void *Arguments, uint32_t) { int32_t omp_vprintf(const char *Format, void *Arguments, uint32_t) {
return vprintf(Format, Arguments); return vprintf(Format, Arguments);
} }
} // namespace impl } // namespace impl
#pragma omp end declare variant #pragma omp end declare variant
Show All 29 Lines

openmp/libomptarget/DeviceRTL/src/Mapping.cpp

Show First 20 Lines • Show All 92 Lines • ▼ Show 20 Lines
#pragma omp end declare variant #pragma omp end declare variant
///} ///}
/// NVPTX Implementation /// NVPTX Implementation
/// ///
///{ ///{
#pragma omp begin declare variant match( \ #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() { uint32_t getNumHardwareThreadsInBlock() {
return __nvvm_read_ptx_sreg_ntid_x(); return __nvvm_read_ptx_sreg_ntid_x();
} }
const llvm::omp::GV &getGridValue() { return llvm::omp::NVPTXGridValues; } const llvm::omp::GV &getGridValue() { return llvm::omp::NVPTXGridValues; }
LaneMaskTy activemask() { LaneMaskTy activemask() {
▲ Show 20 Lines • Show All 181 Lines • Show Last 20 Lines

openmp/libomptarget/DeviceRTL/src/Misc.cpp

Show All 37 Lines
} }
#pragma omp end declare variant #pragma omp end declare variant
/// NVPTX Implementation /// NVPTX Implementation
/// ///
///{ ///{
#pragma omp begin declare variant match( \ #pragma omp begin declare variant match( \
device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)}) device = {arch(nvptx, nvptx64)}, \
implementation = {extension(match_any)})
double getWTick() { double getWTick() {
// Timer precision is 1ns // Timer precision is 1ns
return ((double)1E-9); return ((double)1E-9);
} }
double getWTime() { double getWTime() {
unsigned long long nsecs; unsigned long long nsecs;
Show All 31 Lines

openmp/libomptarget/DeviceRTL/src/Synchronization.cpp

Show First 20 Lines • Show All 285 Lines • ▼ Show 20 Linesvoid unsetCriticalLock(omp_lock_t *Lock) {
(void)atomicExchange((uint32_t *)Lock, UNSET, atomic::acq_rel); (void)atomicExchange((uint32_t *)Lock, UNSET, atomic::acq_rel);
} }
void setCriticalLock(omp_lock_t *Lock) { void setCriticalLock(omp_lock_t *Lock) {
uint64_t LowestActiveThread = utils::ffs(mapping::activemask()) - 1; uint64_t LowestActiveThread = utils::ffs(mapping::activemask()) - 1;
if (mapping::getThreadIdInWarp() == LowestActiveThread) { if (mapping::getThreadIdInWarp() == LowestActiveThread) {
fenceKernel(atomic::release); fenceKernel(atomic::release);
while (!atomicCAS((uint32_t *)Lock, UNSET, SET, atomic::relaxed, while (!atomicCAS((uint32_t *)Lock, UNSET, SET, atomic::relaxed,
atomic::relaxed)) { atomic::relaxed)) {
__builtin_amdgcn_s_sleep(32); __builtin_amdgcn_s_sleep(32);
} }
fenceKernel(atomic::aquire); fenceKernel(atomic::aquire);
} }
} }
#pragma omp end declare variant #pragma omp end declare variant
///} ///}
/// NVPTX Implementation /// NVPTX Implementation
/// ///
///{ ///{
#pragma omp begin declare variant match( \ #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, uint32_t atomicInc(uint32_t *Address, uint32_t Val,
atomic::OrderingTy Ordering) { atomic::OrderingTy Ordering) {
return __nvvm_atom_inc_gen_ui(Address, Val); return __nvvm_atom_inc_gen_ui(Address, Val);
} }
void namedBarrierInit() {} void namedBarrierInit() {}
▲ Show 20 Lines • Show All 161 Lines • ▼ Show 20 Lines
#undef ATOMIC_COMMON_OP #undef ATOMIC_COMMON_OP
#undef ATOMIC_INT_OP #undef ATOMIC_INT_OP
#undef ATOMIC_FP_OP #undef ATOMIC_FP_OP
uint32_t atomic::inc(uint32_t *Addr, uint32_t V, atomic::OrderingTy Ordering) { uint32_t atomic::inc(uint32_t *Addr, uint32_t V, atomic::OrderingTy Ordering) {
return impl::atomicInc(Addr, V, Ordering); return impl::atomicInc(Addr, V, Ordering);
} }
void unsetCriticalLock(omp_lock_t *Lock) { void unsetCriticalLock(omp_lock_t *Lock) { impl::unsetLock(Lock); }
impl::unsetLock(Lock);
}
void setCriticalLock(omp_lock_t *Lock) { void setCriticalLock(omp_lock_t *Lock) { impl::setLock(Lock); }
impl::setLock(Lock);
}
extern "C" { extern "C" {
void __kmpc_ordered(IdentTy *Loc, int32_t TId) { FunctionTracingRAII(); } void __kmpc_ordered(IdentTy *Loc, int32_t TId) { FunctionTracingRAII(); }
void __kmpc_end_ordered(IdentTy *Loc, int32_t TId) { FunctionTracingRAII(); } void __kmpc_end_ordered(IdentTy *Loc, int32_t TId) { FunctionTracingRAII(); }
int32_t __kmpc_cancel_barrier(IdentTy *Loc, int32_t TId) { int32_t __kmpc_cancel_barrier(IdentTy *Loc, int32_t TId) {
FunctionTracingRAII(); FunctionTracingRAII();
▲ Show 20 Lines • Show All 88 Lines • Show Last 20 Lines

openmp/libomptarget/DeviceRTL/src/Utils.cpp

Show First 20 Lines • Show All 53 Lines • ▼ Show 20 Lines
#pragma omp end declare variant #pragma omp end declare variant
///} ///}
/// NVPTX Implementation /// NVPTX Implementation
/// ///
///{ ///{
#pragma omp begin declare variant match( \ #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) { void Unpack(uint64_t Val, uint32_t *LowBits, uint32_t *HighBits) {
uint32_t LowBitsLocal, HighBitsLocal; uint32_t LowBitsLocal, HighBitsLocal;
asm("mov.b64 {%0,%1}, %2;" asm("mov.b64 {%0,%1}, %2;"
: "=r"(LowBitsLocal), "=r"(HighBitsLocal) : "=r"(LowBitsLocal), "=r"(HighBitsLocal)
: "l"(Val)); : "l"(Val));
*LowBits = LowBitsLocal; *LowBits = LowBitsLocal;
*HighBits = HighBitsLocal; *HighBits = HighBitsLocal;
Show All 27 Lines
int32_t shuffleDown(uint64_t Mask, int32_t Var, uint32_t LaneDelta, int32_t shuffleDown(uint64_t Mask, int32_t Var, uint32_t LaneDelta,
int32_t Width) { int32_t Width) {
int Self = mapping::getThreadIdInWarp(); int Self = mapping::getThreadIdInWarp();
int Index = Self + LaneDelta; int Index = Self + LaneDelta;
Index = (int)(LaneDelta + (Self & (Width - 1))) >= Width ? Self : Index; Index = (int)(LaneDelta + (Self & (Width - 1))) >= Width ? Self : Index;
return __builtin_amdgcn_ds_bpermute(Index << 2, Var); return __builtin_amdgcn_ds_bpermute(Index << 2, Var);
} }
bool isSharedMemPtr(const void * Ptr) { bool isSharedMemPtr(const void *Ptr) {
return __builtin_amdgcn_is_shared((const __attribute__((address_space(0))) void *)Ptr); return __builtin_amdgcn_is_shared(
(const __attribute__((address_space(0))) void *)Ptr);
} }
#pragma omp end declare variant #pragma omp end declare variant
///} ///}
/// NVPTX Implementation /// NVPTX Implementation
/// ///
///{ ///{
#pragma omp begin declare variant match( \ #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) { int32_t shuffle(uint64_t Mask, int32_t Var, int32_t SrcLane) {
return __nvvm_shfl_sync_idx_i32(Mask, Var, SrcLane, 0x1f); return __nvvm_shfl_sync_idx_i32(Mask, Var, SrcLane, 0x1f);
} }
int32_t shuffleDown(uint64_t Mask, int32_t Var, uint32_t Delta, int32_t Width) { int32_t shuffleDown(uint64_t Mask, int32_t Var, uint32_t Delta, int32_t Width) {
int32_t T = ((mapping::getWarpSize() - Width) << 8) | 0x1f; int32_t T = ((mapping::getWarpSize() - Width) << 8) | 0x1f;
return __nvvm_shfl_sync_down_i32(Mask, Var, Delta, T); return __nvvm_shfl_sync_down_i32(Mask, Var, Delta, T);
▲ Show 20 Lines • Show All 44 Lines • Show Last 20 Lines

openmp/libomptarget/DeviceRTL/src/Workshare.cpp

Show First 20 Lines • Show All 108 Lines • ▼ Show 20 Lines static void ForStaticNoChunk(int &last, T &lb, T &ub, ST &stride, ST &chunk,
ub = lb + chunk - 1; // Clang uses i <= ub ub = lb + chunk - 1; // Clang uses i <= ub
last = lb <= inputUb && inputUb <= ub; last = lb <= inputUb && inputUb <= ub;
stride = loopSize; // make sure we only do 1 chunk per warp stride = loopSize; // make sure we only do 1 chunk per warp
} }
//////////////////////////////////////////////////////////////////////////////// ////////////////////////////////////////////////////////////////////////////////
// Support for Static Init // Support for Static Init
static void for_static_init(int32_t, int32_t schedtype, static void for_static_init(int32_t, int32_t schedtype, int32_t *plastiter,
int32_t *plastiter, T *plower, T *pupper, T *plower, T *pupper, ST *pstride, ST chunk,
ST *pstride, ST chunk, bool IsSPMDExecutionMode) { bool IsSPMDExecutionMode) {
int32_t gtid = omp_get_thread_num(); int32_t gtid = omp_get_thread_num();
int numberOfActiveOMPThreads = omp_get_num_threads(); int numberOfActiveOMPThreads = omp_get_num_threads();
// All warps that are in excess of the maximum requested, do // All warps that are in excess of the maximum requested, do
// not execute the loop // not execute the loop
ASSERT0(LT_FUSSY, gtid < numberOfActiveOMPThreads, ASSERT0(LT_FUSSY, gtid < numberOfActiveOMPThreads,
"current thread is not needed here; error"); "current thread is not needed here; error");
▲ Show 20 Lines • Show All 536 Lines • Show Last 20 Lines