[all-commits] [llvm/llvm-project] 06bdff: [AMDGPU] Expose llvm fence instruction as clang in...

[Saiyedul Islam via All-commits]

  Branch: refs/heads/master
  Home:   https://github.com/llvm/llvm-project
  Commit: 06bdffb2bb45d8666ec86782d21214ef545a71fd
      https://github.com/llvm/llvm-project/commit/06bdffb2bb45d8666ec86782d21214ef545a71fd
  Author: Saiyedul Islam <Saiyedul.Islam at amd.com>
  Date:   2020-04-27 (Mon, 27 Apr 2020)

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

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