58 lines · plain
1// REQUIRES: amdgpu-registered-target2// REQUIRES: x86-registered-target3 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 %s6 7// AMDGCN has storage-only support for bf16. R600 does not support it should error out when8// 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}}16typedef __attribute__((ext_vector_type(2))) __bf16 bf16_x2;17// r600-error@+1 {{__bf16 is not supported on this target}}18typedef __attribute__((ext_vector_type(4))) __bf16 bf16_x4;19// r600-error@+1 {{__bf16 is not supported on this target}}20typedef __attribute__((ext_vector_type(8))) __bf16 bf16_x8;21// r600-error@+1 {{__bf16 is not supported on this target}}22typedef __attribute__((ext_vector_type(16))) __bf16 bf16_x16;23 24// r600-error@+1 2 {{__bf16 is not supported on this target}}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}}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}}57typedef __bf16 foo __attribute__((__vector_size__(16), __aligned__(16)));58