[AMDGPU] Expose llvm fence instruction as clang intrinsic
authorSaiyedul Islam <Saiyedul.Islam@amd.com>
Mon, 27 Apr 2020 02:56:03 +0000 (08:26 +0530)
committerSameer Sahasrabuddhe <sameer.sahasrabuddhe@amd.com>
Mon, 27 Apr 2020 04:09:03 +0000 (09:39 +0530)
commit06bdffb2bb45d8666ec86782d21214ef545a71fd
treebc5b3122b1f730d6ca2e51bb96da520e2e18ef6a
parent84eff8cef61df89c54e73f99baa3ba3cf03b5e55
[AMDGPU] Expose llvm fence instruction as clang intrinsic

Expose llvm fence instruction as clang builtin for AMDGPU target

__builtin_amdgcn_fence(unsigned int memoryOrdering, const char *syncScope)

The first 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 fence instruction using LLVM atomic C
ABI. The second argument is an AMDGPU-specific synchronization scope
defined as string.

Reviewed By: sameerds

Differential Revision: https://reviews.llvm.org/D75917
clang/include/clang/Basic/BuiltinsAMDGPU.def
clang/include/clang/Sema/Sema.h
clang/lib/CodeGen/CGBuiltin.cpp
clang/lib/Sema/SemaChecking.cpp
clang/test/CodeGenCXX/builtin-amdgcn-fence.cpp [new file with mode: 0644]
clang/test/Sema/builtin-amdgcn-fence-failure.cpp [new file with mode: 0644]
clang/test/SemaOpenCL/builtins-amdgcn-error.cl