OSDN Git Service

[libomptarget] Implement atomic inc and fence functions for AMDGCN using clang builtins
authorSaiyedul Islam <Saiyedul.Islam@amd.com>
Tue, 7 Jul 2020 06:15:26 +0000 (06:15 +0000)
committerSaiyedul Islam <Saiyedul.Islam@amd.com>
Tue, 7 Jul 2020 06:36:25 +0000 (06:36 +0000)
commit38d6640ba55e52a4ae23059164660075c3d8e18a
tree885418e3df9c228b002f9b79642098c791c37eb8
parent0882c9d4fc49858338c9655154f1ad8357a8e516
[libomptarget] Implement atomic inc and fence functions for AMDGCN using clang builtins

This function uses __builtin_amdgcn_atomic_inc32():
  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.

Reviewed By: JonChesterfield

Differential Revision: https://reviews.llvm.org/D83132
openmp/libomptarget/deviceRTLs/amdgcn/src/hip_atomics.h
openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h