clang/test/CodeGenOpenCL/builtins-amdgcn-fp-atomics-gfx11-err.cl

   1 // RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx1100 \
   2 // RUN:   -verify -S -o - %s
   3
   4 // REQUIRES: amdgpu-registered-target
   5
   6 typedef half  __attribute__((ext_vector_type(2))) half2;
   7 typedef short __attribute__((ext_vector_type(2))) short2;
   8
   9
  10 void test_atomic_fadd(__global half2 *addrh2, __local half2 *addrh2l, half2 xh2,
  11                       __global short2 *addrs2, __local short2 *addrs2l, short2 xs2,
  12                       __global float *addrf, float xf) {
  13   __builtin_amdgcn_flat_atomic_fadd_v2f16(addrh2, xh2); // expected-error{{'__builtin_amdgcn_flat_atomic_fadd_v2f16' needs target feature atomic-flat-pk-add-16-insts}}
  14   __builtin_amdgcn_flat_atomic_fadd_v2bf16(addrs2, xs2); // expected-error{{'__builtin_amdgcn_flat_atomic_fadd_v2bf16' needs target feature atomic-flat-pk-add-16-insts}}
  15   __builtin_amdgcn_global_atomic_fadd_v2bf16(addrs2, xs2); // expected-error{{'__builtin_amdgcn_global_atomic_fadd_v2bf16' needs target feature atomic-global-pk-add-bf16-inst}}
  16   __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}}
  17   __builtin_amdgcn_ds_atomic_fadd_v2bf16(addrs2l, xs2); // expected-error{{'__builtin_amdgcn_ds_atomic_fadd_v2bf16' needs target feature atomic-ds-pk-add-16-insts}}
  18   __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}}
  19 }