1 // REQUIRES
: amdgpu-registered-target
2 // RUN
: %clang_cc1 -O0 -cl-std
=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx908 \
3 // RUN
: -verify -S -o - %s
5 typedef half __attribute__
((ext_vector_type(2))) half2
;
7 void test_global_fadd
(__global half2
*addrh2
, __local half2
*addrh2l
, half2 xh2
,
8 __global float
*addrf
, float xf
,
9 __global double
*addr
, double x
) {
13 *half_rtn
= __builtin_amdgcn_global_atomic_fadd_v2f16
(addrh2, xh2
); // expected-error{{'__builtin_amdgcn_global_atomic_fadd_v2f16' needs target feature atomic-buffer-global-pk-add-f16-insts}}
14 *fp_rtn
= __builtin_amdgcn_global_atomic_fadd_f32
(addr, x
); // expected-error{{'__builtin_amdgcn_global_atomic_fadd_f32' needs target feature atomic-fadd-rtn-insts}}
15 *rtn
= __builtin_amdgcn_global_atomic_fadd_f64
(addr, x
); // expected-error{{'__builtin_amdgcn_global_atomic_fadd_f64' needs target feature gfx90a-insts}}
16 *rtn
= __builtin_amdgcn_global_atomic_fmax_f64
(addr, x
); // expected-error{{'__builtin_amdgcn_global_atomic_fmax_f64' needs target feature gfx90a-insts}}
17 *rtn
= __builtin_amdgcn_global_atomic_fmin_f64
(addr, x
); // expected-error{{'__builtin_amdgcn_global_atomic_fmin_f64' needs target feature gfx90a-insts}}
18 *rtn
= __builtin_amdgcn_flat_atomic_fadd_f64
(addr, x
); // expected-error{{'__builtin_amdgcn_flat_atomic_fadd_f64' needs target feature gfx90a-insts}}
19 *rtn
= __builtin_amdgcn_flat_atomic_fmin_f64
(addr, x
); // expected-error{{'__builtin_amdgcn_flat_atomic_fmin_f64' needs target feature gfx90a-insts}}
20 *rtn
= __builtin_amdgcn_flat_atomic_fmax_f64
(addr, x
); // expected-error{{'__builtin_amdgcn_flat_atomic_fmax_f64' needs target feature gfx90a-insts}}
21 __builtin_amdgcn_ds_atomic_fadd_v2f16
(addrh2l, xh2
); // expected-error{{'__builtin_amdgcn_ds_atomic_fadd_v2f16' needs target feature atomic-ds-pk-add-16-insts}}