HomePhabricator

[AMDGPU] Expose llvm fence instruction as clang intrinsic

Authored by saiislam on Apr 26 2020, 7:56 PM.

Description

[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

Details

Committed
sameerdsApr 26 2020, 9:09 PM
Reviewer
sameerds
Differential Revision
D75917: Expose llvm fence instruction as clang intrinsic
Parents
rG84eff8cef61d: [llvm-objcopy][MachO] Fix segment's vmsize
Branches
Unknown
Tags
Unknown