llvm-project/clang/test/Sema/builtin-amdgcn-fence-failure.cpp

// REQUIRES: amdgpu-registered-target
// RUN: not %clang_cc1 %s -o - -S -triple=amdgcn-amd-amdhsa 2>&1 | FileCheck %s

void test_amdgcn_fence_failure() {

  // CHECK: error: Unsupported atomic synchronization scope
  __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "foobar");
}
[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 08:26:03 +05:30			`// REQUIRES: amdgpu-registered-target`
Explicitly tell Clang to output to stdout in a test that runs FileCheck 2020-04-27 10:03:45 +02: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 08:26:03 +05:30
			`void test_amdgcn_fence_failure() {`

			`// CHECK: error: Unsupported atomic synchronization scope`
			`__builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "foobar");`
			`}`