mirror of
https://github.com/llvm/llvm-project.git
synced 2025-05-10 05:36:07 +00:00
06bdffb2bb
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
23 lines
682 B
C++
23 lines
682 B
C++
// REQUIRES: amdgpu-registered-target
|
|
// RUN: %clang_cc1 %s -emit-llvm -O0 -o - \
|
|
// RUN: -triple=amdgcn-amd-amdhsa | opt -S | FileCheck %s
|
|
|
|
void test_memory_fence_success() {
|
|
// CHECK-LABEL: test_memory_fence_success
|
|
|
|
// CHECK: fence syncscope("workgroup") seq_cst
|
|
__builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "workgroup");
|
|
|
|
// CHECK: fence syncscope("agent") acquire
|
|
__builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "agent");
|
|
|
|
// CHECK: fence seq_cst
|
|
__builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "");
|
|
|
|
// CHECK: fence syncscope("agent") acq_rel
|
|
__builtin_amdgcn_fence(4, "agent");
|
|
|
|
// CHECK: fence syncscope("workgroup") release
|
|
__builtin_amdgcn_fence(3, "workgroup");
|
|
}
|