1 // RUN
: %clang_cc1 -O0 -cl-std
=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx90a \
2 // RUN
: -verify -S -o - %s
4 // REQUIRES
: amdgpu-registered-target
6 typedef half __attribute__
((ext_vector_type(2))) half2
;
7 typedef short __attribute__
((ext_vector_type(2))) short2
;
9 void test_atomic_fadd
(__global half2
*addrh2
, half2 xh2
,
10 __global short2
*addrs2
, __local short2
*addrs2l
, short2 xs2
,
11 __global float
*addrf
, float xf
) {
12 __builtin_amdgcn_flat_atomic_fadd_f32
(addrf, xf
); // expected-error{{'__builtin_amdgcn_flat_atomic_fadd_f32' needs target feature gfx940-insts}}
13 __builtin_amdgcn_flat_atomic_fadd_v2f16
(addrh2, xh2
); // expected-error{{'__builtin_amdgcn_flat_atomic_fadd_v2f16' needs target feature gfx940-insts}}
14 __builtin_amdgcn_flat_atomic_fadd_v2bf16
(addrs2, xs2
); // expected-error{{'__builtin_amdgcn_flat_atomic_fadd_v2bf16' needs target feature gfx940-insts}}
15 __builtin_amdgcn_global_atomic_fadd_v2bf16
(addrs2, xs2
); // expected-error{{'__builtin_amdgcn_global_atomic_fadd_v2bf16' needs target feature gfx940-insts}}
16 __builtin_amdgcn_ds_atomic_fadd_v2bf16
(addrs2l, xs2
); // expected-error{{'__builtin_amdgcn_ds_atomic_fadd_v2bf16' needs target feature gfx940-insts}}