xref: /llvm-project/clang/test/Sema/builtin-amdgcn-atomic-inc-dec-failure.cpp (revision 9275e14379961a4304de559f16fdbac275fb6301)
1 // REQUIRES: amdgpu-registered-target
2 // RUN: %clang_cc1 %s -x hip -fcuda-is-device -o - \
3 // RUN:   -triple=amdgcn-amd-amdhsa -fsyntax-only \
4 // RUN:   -verify=dev
5 // RUN: %clang_cc1 %s -x hip -triple x86_64 -o - \
6 // RUN:   -aux-triple amdgcn-amd-amdhsa -fsyntax-only \
7 // RUN:   -verify=host
8 
9 // dev-no-diagnostics
10 
test_host()11 void test_host() {
12   __UINT32_TYPE__ val32;
13   __UINT64_TYPE__ val64;
14 
15   // host-error@+1 {{reference to __device__ function '__builtin_amdgcn_atomic_inc32' in __host__ function}}
16   val32 = __builtin_amdgcn_atomic_inc32(&val32, val32, __ATOMIC_SEQ_CST, "");
17 
18   // host-error@+1 {{reference to __device__ function '__builtin_amdgcn_atomic_inc64' in __host__ function}}
19   val64 = __builtin_amdgcn_atomic_inc64(&val64, val64, __ATOMIC_SEQ_CST, "");
20 
21   // host-error@+1 {{reference to __device__ function '__builtin_amdgcn_atomic_dec32' in __host__ function}}
22   val32 = __builtin_amdgcn_atomic_dec32(&val32, val32, __ATOMIC_SEQ_CST, "");
23 
24   // host-error@+1 {{reference to __device__ function '__builtin_amdgcn_atomic_dec64' in __host__ function}}
25   val64 = __builtin_amdgcn_atomic_dec64(&val64, val64, __ATOMIC_SEQ_CST, "");
26 }
27