[all-commits] [llvm/llvm-project] 0882c9: [AMDGPU] Change Clang AMDGCN atomic inc/dec builti...

[Saiyedul Islam via All-commits]

  Branch: refs/heads/master
  Home:   https://github.com/llvm/llvm-project
  Commit: 0882c9d4fc49858338c9655154f1ad8357a8e516
      https://github.com/llvm/llvm-project/commit/0882c9d4fc49858338c9655154f1ad8357a8e516
  Author: Saiyedul Islam <Saiyedul.Islam at amd.com>
  Date:   2020-07-07 (Tue, 07 Jul 2020)

  Changed paths:
    M clang/include/clang/Basic/BuiltinsAMDGPU.def
    M clang/test/CodeGenCXX/builtin-amdgcn-atomic-inc-dec.cpp
    M clang/test/Sema/builtin-amdgcn-atomic-inc-dec-failure.cpp
    M clang/test/SemaOpenCL/builtins-amdgcn-error.cl

  Log Message:
  -----------
  [AMDGPU] Change Clang AMDGCN atomic inc/dec builtins to take unsigned values

builtin_amdgcn_atomic_inc32(uint *Ptr, uint Val, unsigned MemoryOrdering, const char *SyncScope)
builtin_amdgcn_atomic_inc64(uint64_t *Ptr, uint64_t Val, unsigned MemoryOrdering, const char *SyncScope)
builtin_amdgcn_atomic_dec32(uint *Ptr, uint Val, unsigned MemoryOrdering, const char *SyncScope)
builtin_amdgcn_atomic_dec64(uint64_t *Ptr, uint64_t Val, unsigned MemoryOrdering, const char *SyncScope)

As AMDGCN IR instrinsic for atomic inc/dec does unsigned comparison,
these clang builtins should also take unsigned types instead of signed
int types.

Reviewed By: JonChesterfield

Differential Revision: https://reviews.llvm.org/D83121


  Commit: 38d6640ba55e52a4ae23059164660075c3d8e18a
      https://github.com/llvm/llvm-project/commit/38d6640ba55e52a4ae23059164660075c3d8e18a
  Author: Saiyedul Islam <Saiyedul.Islam at amd.com>
  Date:   2020-07-07 (Tue, 07 Jul 2020)

  Changed paths:
    M openmp/libomptarget/deviceRTLs/amdgcn/src/hip_atomics.h
    M openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h

  Log Message:
  -----------
  [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


Compare: https://github.com/llvm/llvm-project/compare/16f3d698f2af...38d6640ba55e