[PATCH] D80804: [AMDGPU] Expose llvm atomic inc/dec instructions as clang builtins for AMDGPU target

[Saiyedul Islam via Phabricator via cfe-commits]

saiislam created this revision.
saiislam added reviewers: arsenm, sameerds, JonChesterfield, jdoerfert.
Herald added subscribers: cfe-commits, kerbowa, jfb, t-tye, tpr, dstuttard, yaxunl, nhaehnle, wdng, jvesely, kzhuravl.
Herald added a project: clang.
saiislam added parent revisions: D75917: Expose llvm fence instruction as clang intrinsic, D73076: [libomptarget] Implement most hip atomic functions in terms of intrinsics.

__builtin_amdgcn_atomic_inc(int *Ptr, int Val, unsigned MemoryOrdering,

  				const char *SyncScope, bool IsVolatile)

__builtin_amdgcn_atomic_dec(int *Ptr, int Val, unsigned MemoryOrdering,

  const char *SyncScope, bool IsVolatile)

First, second, and fifth argument gets transparently passed to the llvm
intruction. 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 LLVM
atomic C ABI. The fourth argument is an AMDGPU-specific synchronization scope
defined as string.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D80804

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

-------------- next part --------------
A non-text attachment was scrubbed...
Name: D80804.267240.patch
Type: text/x-patch
Size: 19929 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20200529/9bf274d2/attachment-0001.bin>