maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity
llvm-project/clang/test/Sema/builtin-amdgcn-fence-failur...

Ignoring revisions in .git-blame-ignore-revs. Click here to bypass and see the normal blame view.

9 lines
271 B
C++
Raw Normal View History

[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
2020-04-27 10:56:03 +08:00
// REQUIRES: amdgpu-registered-target
Explicitly tell Clang to output to stdout in a test that runs FileCheck
2020-04-27 16:03:45 +08:00
// RUN: not %clang_cc1 %s -o - -S -triple=amdgcn-amd-amdhsa 2>&1 | FileCheck %s
[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
2020-04-27 10:56:03 +08:00
void test_amdgcn_fence_failure() {
// CHECK: error: Unsupported atomic synchronization scope
__builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "foobar");
}