[Openmp-commits] [PATCH] D83132: [libomptarget] Implement atomic inc and fence functions for AMDGCN using clang builtins

[Jon Chesterfield via Phabricator via Openmp-commits]

JonChesterfield added inline comments.


================
Comment at: openmp/libomptarget/deviceRTLs/amdgcn/src/amdgcn_threadfence.hip:11
+
+EXTERN void __kmpc_impl_threadfence() {
+  __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "agent");
----------------
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.


================
Comment at: openmp/libomptarget/deviceRTLs/amdgcn/src/hip_atomics.h:37
+// Variants of atomicInc with and without wraparound MAX value
+DEVICE uint32_t atomicInc(uint32_t *address) {
+  return __builtin_amdgcn_atomic_inc32(address, UINT32_MAX, __ATOMIC_SEQ_CST,
----------------
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.


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D83132/new/

https://reviews.llvm.org/D83132