1 // REQUIRES: amdgpu-registered-target
2 // REQUIRES: x86-registered-target
3
4 // RUN: %clang_cc1 "-aux-triple" "x86_64-unknown-linux-gnu" "-triple" "r600-unknown-unknown"\
5 // RUN: -fcuda-is-device "-aux-target-cpu" "x86-64" -fsyntax-only -verify=r600 %s
6
7 // AMDGCN has storage-only support for bf16. R600 does not support it should error out when
8 // it's the main target.
9
10 #include "Inputs/cuda.h"
11
12 // There should be no errors on using the type itself, or when loading/storing values for amdgcn.
13 // r600 should error on all uses of the type.
14
15 // r600-error@+1 {{__bf16 is not supported on this target}}
16 typedef __attribute__((ext_vector_type(2))) __bf16 bf16_x2;
17 // r600-error@+1 {{__bf16 is not supported on this target}}
18 typedef __attribute__((ext_vector_type(4))) __bf16 bf16_x4;
19 // r600-error@+1 {{__bf16 is not supported on this target}}
20 typedef __attribute__((ext_vector_type(8))) __bf16 bf16_x8;
21 // r600-error@+1 {{__bf16 is not supported on this target}}
22 typedef __attribute__((ext_vector_type(16))) __bf16 bf16_x16;
23
24 // r600-error@+1 2 {{__bf16 is not supported on this target}}
test(bool b,__bf16 * out,__bf16 in)25 __device__ void test(bool b, __bf16 *out, __bf16 in) {
26 __bf16 bf16 = in; // r600-error {{__bf16 is not supported on this target}}
27 *out = bf16;
28
29 // r600-error@+1 {{__bf16 is not supported on this target}}
30 typedef __attribute__((ext_vector_type(2))) __bf16 bf16_x2;
31 bf16_x2 vec2_a, vec2_b;
32 vec2_a = vec2_b;
33
34 // r600-error@+1 {{__bf16 is not supported on this target}}
35 typedef __attribute__((ext_vector_type(4))) __bf16 bf16_x4;
36 bf16_x4 vec4_a, vec4_b;
37 vec4_a = vec4_b;
38
39 // r600-error@+1 {{__bf16 is not supported on this target}}
40 typedef __attribute__((ext_vector_type(8))) __bf16 bf16_x8;
41 bf16_x8 vec8_a, vec8_b;
42 vec8_a = vec8_b;
43
44 // r600-error@+1 {{__bf16 is not supported on this target}}
45 typedef __attribute__((ext_vector_type(16))) __bf16 bf16_x16;
46 bf16_x16 vec16_a, vec16_b;
47 vec16_a = vec16_b;
48 }
49
50 // r600-error@+1 2 {{__bf16 is not supported on this target}}
hostfn(__bf16 a)51 __bf16 hostfn(__bf16 a) {
52 return a;
53 }
54
55 // r600-error@+2 {{__bf16 is not supported on this target}}
56 // r600-error@+1 {{vector size not an integral multiple of component size}}
57 typedef __bf16 foo __attribute__((__vector_size__(16), __aligned__(16)));
58