1 // REQUIRES: amdgpu-registered-target
2 // REQUIRES: x86-registered-target
4 // RUN: %clang_cc1 "-triple" "x86_64-unknown-linux-gnu" "-aux-triple" "amdgcn-amd-amdhsa"\
5 // RUN: "-target-cpu" "x86-64" -fsyntax-only -verify=amdgcn %s
6 // RUN: %clang_cc1 "-aux-triple" "x86_64-unknown-linux-gnu" "-triple" "amdgcn-amd-amdhsa"\
7 // RUN: -fcuda-is-device "-aux-target-cpu" "x86-64" -fsyntax-only -verify=amdgcn %s
9 // RUN: %clang_cc1 "-aux-triple" "x86_64-unknown-linux-gnu" "-triple" "r600-unknown-unknown"\
10 // RUN: -fcuda-is-device "-aux-target-cpu" "x86-64" -fsyntax-only -verify=amdgcn,r600 %s
12 // AMDGCN has storage-only support for bf16. R600 does not support it should error out when
13 // it's the main target.
15 #include "Inputs/cuda.h"
17 // There should be no errors on using the type itself, or when loading/storing values for amdgcn.
18 // r600 should error on all uses of the type.
20 // r600-error@+1 {{__bf16 is not supported on this target}}
21 typedef __attribute__((ext_vector_type(2))) __bf16 bf16_x2;
22 // r600-error@+1 {{__bf16 is not supported on this target}}
23 typedef __attribute__((ext_vector_type(4))) __bf16 bf16_x4;
24 // r600-error@+1 {{__bf16 is not supported on this target}}
25 typedef __attribute__((ext_vector_type(8))) __bf16 bf16_x8;
26 // r600-error@+1 {{__bf16 is not supported on this target}}
27 typedef __attribute__((ext_vector_type(16))) __bf16 bf16_x16;
29 // r600-error@+1 2 {{__bf16 is not supported on this target}}
30 __device__ void test(bool b, __bf16 *out, __bf16 in) {
31 __bf16 bf16 = in; // r600-error {{__bf16 is not supported on this target}}
33 bf16 + bf16; // amdgcn-error {{invalid operands to binary expression ('__bf16' and '__bf16')}}
34 bf16 - bf16; // amdgcn-error {{invalid operands to binary expression ('__bf16' and '__bf16')}}
35 bf16 * bf16; // amdgcn-error {{invalid operands to binary expression ('__bf16' and '__bf16')}}
36 bf16 / bf16; // amdgcn-error {{invalid operands to binary expression ('__bf16' and '__bf16')}}
40 bf16 + fp16; // amdgcn-error {{invalid operands to binary expression ('__bf16' and '__fp16')}}
41 fp16 + bf16; // amdgcn-error {{invalid operands to binary expression ('__fp16' and '__bf16')}}
42 bf16 - fp16; // amdgcn-error {{invalid operands to binary expression ('__bf16' and '__fp16')}}
43 fp16 - bf16; // amdgcn-error {{invalid operands to binary expression ('__fp16' and '__bf16')}}
44 bf16 * fp16; // amdgcn-error {{invalid operands to binary expression ('__bf16' and '__fp16')}}
45 fp16 * bf16; // amdgcn-error {{invalid operands to binary expression ('__fp16' and '__bf16')}}
46 bf16 / fp16; // amdgcn-error {{invalid operands to binary expression ('__bf16' and '__fp16')}}
47 fp16 / bf16; // amdgcn-error {{invalid operands to binary expression ('__fp16' and '__bf16')}}
48 bf16 = fp16; // amdgcn-error {{assigning to '__bf16' from incompatible type '__fp16'}}
49 fp16 = bf16; // amdgcn-error {{assigning to '__fp16' from incompatible type '__bf16'}}
50 bf16 + (b ? fp16 : bf16); // amdgcn-error {{incompatible operand types ('__fp16' and '__bf16')}}
53 // amdgcn-error@+1 {{static_cast from '__bf16' to 'unsigned short' is not allowed}}
54 unsigned short u16bf16 = static_cast<unsigned short>(bf16);
55 // amdgcn-error@+2 {{C-style cast from 'unsigned short' to '__bf16' is not allowed}}
56 // r600-error@+1 {{__bf16 is not supported on this target}}
57 bf16 = (__bf16)u16bf16;
59 // amdgcn-error@+1 {{static_cast from '__bf16' to 'float' is not allowed}}
60 float f32bf16 = static_cast<float>(bf16);
61 // amdgcn-error@+2 {{C-style cast from 'float' to '__bf16' is not allowed}}
62 // r600-error@+1 {{__bf16 is not supported on this target}}
63 bf16 = (__bf16)f32bf16;
65 // amdgcn-error@+1 {{static_cast from '__bf16' to 'double' is not allowed}}
66 double f64bf16 = static_cast<double>(bf16);
67 // amdgcn-error@+2 {{C-style cast from 'double' to '__bf16' is not allowed}}
68 // r600-error@+1 {{__bf16 is not supported on this target}}
69 bf16 = (__bf16)f64bf16;
71 // r600-error@+1 {{__bf16 is not supported on this target}}
72 typedef __attribute__((ext_vector_type(2))) __bf16 bf16_x2;
73 bf16_x2 vec2_a, vec2_b;
76 // r600-error@+1 {{__bf16 is not supported on this target}}
77 typedef __attribute__((ext_vector_type(4))) __bf16 bf16_x4;
78 bf16_x4 vec4_a, vec4_b;
81 // r600-error@+1 {{__bf16 is not supported on this target}}
82 typedef __attribute__((ext_vector_type(8))) __bf16 bf16_x8;
83 bf16_x8 vec8_a, vec8_b;
86 // r600-error@+1 {{__bf16 is not supported on this target}}
87 typedef __attribute__((ext_vector_type(16))) __bf16 bf16_x16;
88 bf16_x16 vec16_a, vec16_b;
92 // r600-error@+1 2 {{__bf16 is not supported on this target}}
93 __bf16 hostfn(__bf16 a) {
97 // r600-error@+2 {{__bf16 is not supported on this target}}
98 // r600-error@+1 {{vector size not an integral multiple of component size}}
99 typedef __bf16 foo __attribute__((__vector_size__(16), __aligned__(16)));