[all-commits] [llvm/llvm-project] 675cef: [AMDGPU] Introduce Clang builtins to be mapped to ...

[Saiyedul Islam via All-commits]

  Branch: refs/heads/master
  Home:   https://github.com/llvm/llvm-project
  Commit: 675cefbf60270f59057972e33365a09590fb3694
      https://github.com/llvm/llvm-project/commit/675cefbf60270f59057972e33365a09590fb3694
  Author: Saiyedul Islam <Saiyedul.Islam at amd.com>
  Date:   2020-06-09 (Tue, 09 Jun 2020)

  Changed paths:
    M clang/include/clang/Basic/BuiltinsAMDGPU.def
    M clang/lib/CodeGen/CGBuiltin.cpp
    M clang/lib/CodeGen/CodeGenFunction.h
    M clang/lib/Sema/SemaChecking.cpp
    A clang/test/CodeGenCXX/builtin-amdgcn-atomic-inc-dec.cpp
    A clang/test/Sema/builtin-amdgcn-atomic-inc-dec-failure.cpp
    M clang/test/SemaOpenCL/builtins-amdgcn-error.cl

  Log Message:
  -----------
  [AMDGPU] Introduce Clang builtins to be mapped to AMDGCN atomic inc/dec intrinsics

Summary:
__builtin_amdgcn_atomic_inc32(int *Ptr, int Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_inc64(int64_t *Ptr, int64_t Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_dec32(int *Ptr, int Val, unsigned MemoryOrdering, const char *SyncScope)
__builtin_amdgcn_atomic_dec64(int64_t *Ptr, int64_t Val, unsigned MemoryOrdering, const char *SyncScope)

First and second arguments gets transparently passed to the amdgcn atomic
inc/dec intrinsic. Fifth argument of the intrinsic is set as true if the
first argument of the builtin is a volatile pointer. The third argument of
this builtin is one of the memory-ordering specifiers ATOMIC_ACQUIRE,
ATOMIC_RELEASE, ATOMIC_ACQ_REL, or ATOMIC_SEQ_CST following C++11 memory
model semantics. This is mapped to corresponding LLVM atomic memory ordering
for the atomic inc/dec instruction using CLANG atomic C ABI. The fourth
argument is an AMDGPU-specific synchronization scope defined as string.

Reviewers: arsenm, sameerds, JonChesterfield, jdoerfert

Reviewed By: arsenm, sameerds

Subscribers: kzhuravl, jvesely, wdng, nhaehnle, yaxunl, dstuttard, tpr, t-tye, jfb, kerbowa, cfe-commits

Tags: #clang

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