This is an archive of the discontinued LLVM Phabricator instance.

[libomptarget] Implement atomic inc and fence functions for AMDGCN using clang builtins
ClosedPublic

Authored by saiislam on Jul 3 2020, 7:41 AM.

Details

Summary

These functions use __builtin_amdgcn_atomic_inc32():

uint32_t atomicInc(uint32_t *address);
uint32_t atomicInc(uint32_t *address, uint32_t max);

These functions use builtin_amdgcn_fence():
kmpc_impl_threadfence()
kmpc_impl_threadfence_block()
kmpc_impl_threadfence_system()

They will take place of current mechanism of directly calling IR functions.

Diff Detail

Event Timeline

saiislam created this revision.Jul 3 2020, 7:41 AM
Herald added a project: Restricted Project. · View Herald TranscriptJul 3 2020, 7:41 AM
JonChesterfield added inline comments.Jul 3 2020, 8:10 AM
openmp/libomptarget/deviceRTLs/amdgcn/src/amdgcn_threadfence.hip
11

Since these are one line wrappers around intrinsics, it's probably better to implement them as INLINE annotated functions in target_impl.h. Less noise in the filesystem, can do some optimisation on the single-tu level before calling opt.

openmp/libomptarget/deviceRTLs/amdgcn/src/hip_atomics.h
37

Please delete the unused one. Also, should need to mark it INLINE instead of DEVICE to avoid duplicate symbol errors. I'm surprised this linked.

saiislam updated this revision to Diff 275422.Jul 3 2020, 9:56 AM
  1. Removed unused functions
  2. Changed to INLINE
  3. Moved fence functions to target_impl.h
JonChesterfield accepted this revision.Jul 3 2020, 10:31 AM

This looks right, thank you. We need to build with trunk clang at present anyway. I think we'll be able to build with 11.0 once that forks.

This revision is now accepted and ready to land.Jul 3 2020, 10:31 AM
This revision was automatically updated to reflect the committed changes.