In D80804#2063451, @sameerds wrote:

Actually, the question really is about why inc/dec are needed as separate operations either as IR intrinsics or Clang builtins. Why not just expose a __builtin_amdgcn_atomicrmw that takes a scope, and map it to the LLVM atomicrmw? That would be way cleaner. The language can provide convenience functions for inc/dec that internally call the rmw builtin.

1. Updated title and description.
2. Replaced pointer element type usage to i32 type for getIntrinsic.
3. Volatile argument of the instrinsic now comes from pointer type of the first argument of the builtin.
4. Updated all test cases wrt above change.
5. Added test case for volatile pointer type.

Shifted test cases in openmp-offload-gpu.c for better visual segmentation.
Updated device function test case to be more accurate.

Added test case to show treatment of specific functions as builtins or functions on the device

Fixed typo. Added test case to check for device side functions.

Removed isOpenMPGPU() to avoid defining OpenMP compatibility of an architecture.
Reverting back to explicitly checking NVPTX and AMDGCN architectures. Also, split
handling of NVPTX's and AMDGCN's handling of getBuiltinID. For AMDGCN it now uses
OpenMPIsDevice LangOpt and returns 0 for every device library function, except for
printf and malloc.

Moved isGPU() from llvm's Triple.h to clang's TargetInfo. Renamed it to isOpenMPGPU()
to represent target's compatibility with OpenMP offloading and reduce its scope.

Added test cases. Added a wrapper isGPU() for isNVPTX()/isAMDGCN().

Fixed lint errors.

Removed a break

Switched the nested loops in lowerShiftReservedVGPR() to move
sorting of unique live-ins to the outer loop. Improved cleanup of
reserved vgpr after lowering it to a lower value, and when it needs
to be removed.

Handled a test case which was failing in debug only.

Removed unsued frame index argument from vgpr reserving function. Added condition to reserve vgpr only if machine function has stack objects.

Cleaned code for reserving vgpr.

Fixed failing test cases. Now VGPR reservation happens in an
independent function and doesn't overload allocateSGPRSpillToVGPR.
Updated test case to hard code usage of all but one VGPRs and use
it for SGPR spill.

Handled few seg faults in the lit-tests.

Removed commented code.

Rebased and added code for proper lower shifting of reserved VGPR.

Updated description and added a failing test case for integer scope.

Added check and test for sync scope to be a constant literal.

Removed documentation from clang doc. Squashed all changes into a single commit.

Now highest available VGPR is reserved in the beginning and it is
switched with lowest available VGPR after RA.

Moved builtin-amdgcn-fence-failure.cpp from clang/test/CodeGenCXX/
to clang/test/Sema/ since it is checking an error.

Removed stale commented code

Changed the builtin to be AMDGCN-specific

1. Moved test case to clang/test/CodeGenCXX.
2. Added a failing test case with invalid sync scope, which gets detected by implementation of fence instruction.
3. Updated the change description of the builtin.
4. Updated the clang documentation describing mapping of C++ memory-ordering to LLVM memory-ordering.

Removed reserved VGPR from RA list. Added a test case which
ensures spilling of SGPR into reserved VGPR.

Changes:

1. Moved builtin definition with rest of atomic builtins
2. Updated validity checking of memory order using exact mathches instead of range matching
3. Added a sucessful test case which passes arbitrary string scope
4. Corrected formatting

Removed OpenCL specific dependencies

Corrected formatting