This is an archive of the discontinued LLVM Phabricator instance.

[libomptarget][amdgcn] Build amdgcn devicertl as openmp
ClosedPublic

Authored by JonChesterfield on Feb 11 2021, 11:48 AM.

Details

Summary

[libomptarget][amdgcn] Build amdgcn devicertl as openmp

Change cmake to build as openmp and fix up some minor errors in the code.

Diff Detail

Event Timeline

JonChesterfield requested review of this revision.Feb 11 2021, 11:48 AM
Herald added a project: Restricted Project. · View Herald TranscriptFeb 11 2021, 11:48 AM

This doesn't toggle the cmake to build the amdgcn plugin by default. I'm planning to do that after the dust settles in the internal builds from this.

This doesn't toggle the cmake to build the amdgcn plugin by default. I'm planning to do that after the dust settles in the internal builds from this.

One nice thing if AMD can build deviceRTLs with OpenMP is, we can completely get rid of the macro DEVICE.

JonChesterfield added a comment.EditedFeb 11 2021, 11:57 AM

Yep. I'm walking a slightly fine line between this and the internal rocm builds so would like to keep the functional changes separate from the cleanup. Plan is ship this, which 'should' work internally, pick up any pieces if it doesn't, then drop the DEVICE and probably INLINE macros. Maybe get rid of some of EXTERN as well.

Curiously the internal build does compile this library as openmp, and has done since December, but uses a slightly different clang and a slightly different devicertl.

jdoerfert added inline comments.Feb 11 2021, 12:01 PM
openmp/libomptarget/deviceRTLs/amdgcn/src/amdgcn_intrinsics.h
46 ↗(On Diff #323105)

Do we need to declare builtins? I guess we can but we should not have to, right?

openmp/libomptarget/deviceRTLs/amdgcn/src/amdgcn_intrinsics.h
46 ↗(On Diff #323105)

Today, yes. I think the 'right' behaviour is to allow these within an appropriate declare variant, and an adequate approximation would be to allow them everywhere.

As raw openmp, the assumption is that any function can be called from the host. That doesn't mean anything for __builtin_amdgcn_dispatch_ptr, so it's reasonable that these aren't automatically available.

At present though, I'm just trying to get this library to compile with trunk clang as-is. Curiously, if one of these is declared with the wrong type, clang emits an error on the declaration.

jdoerfert added inline comments.Feb 11 2021, 1:28 PM
openmp/libomptarget/deviceRTLs/amdgcn/src/amdgcn_intrinsics.h
46 ↗(On Diff #323105)

I can just call these like other builtins (when I target amdgcn):
https://godbolt.org/z/nabGT7

  • drop unused declaration
openmp/libomptarget/deviceRTLs/amdgcn/src/amdgcn_intrinsics.h
46 ↗(On Diff #323105)

That is interesting.

I wrote this patch against a version of --cuda-device-only. These may have only raised an error on the host pass, which is not invoked by the present -Xclang themed cmake.

  • Drop amdgcn intrinsics, not needed with host pass disabled
This revision is now accepted and ready to land.Feb 11 2021, 5:42 PM
JonChesterfield edited the summary of this revision. (Show Details)Feb 12 2021, 1:50 AM
This revision was landed with ongoing or failed builds.Feb 12 2021, 1:51 AM
This revision was automatically updated to reflect the committed changes.