1 // REQUIRES: amdgpu-registered-target
2 // RUN: not %clang_cc1 %s -x hip -fcuda-is-device -o - -emit-llvm -triple=amdgcn-amd-amdhsa 2>&1 | FileCheck %s
3
test_host()4 void test_host() {
5 __UINT32_TYPE__ val32;
6 __UINT64_TYPE__ val64;
7
8 // CHECK: error: reference to __device__ function '__builtin_amdgcn_atomic_inc32' in __host__ function
9 val32 = __builtin_amdgcn_atomic_inc32(&val32, val32, __ATOMIC_SEQ_CST, "");
10
11 // CHECK: error: reference to __device__ function '__builtin_amdgcn_atomic_inc64' in __host__ function
12 val64 = __builtin_amdgcn_atomic_inc64(&val64, val64, __ATOMIC_SEQ_CST, "");
13
14 // CHECK: error: reference to __device__ function '__builtin_amdgcn_atomic_dec32' in __host__ function
15 val32 = __builtin_amdgcn_atomic_dec32(&val32, val32, __ATOMIC_SEQ_CST, "");
16
17 // CHECK: error: reference to __device__ function '__builtin_amdgcn_atomic_dec64' in __host__ function
18 val64 = __builtin_amdgcn_atomic_dec64(&val64, val64, __ATOMIC_SEQ_CST, "");
19 }
20