| // Test target codegen - host bc file has to be created first. |
| // RUN: %clang_cc1 -fopenmp -x c++ -triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-host.bc |
| // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -fsyntax-only |
| // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-linux-gnu -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-host.bc |
| // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-linux-gnu -fopenmp-targets=nvptx64-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -fsyntax-only |
| |
| struct T { |
| char a; |
| #ifndef _ARCH_PPC |
| // expected-note@+1 {{'f' defined here}} |
| __float128 f; |
| #else |
| // expected-note@+1 {{'f' defined here}} |
| long double f; |
| #endif |
| char c; |
| T() : a(12), f(15) {} |
| #ifndef _ARCH_PPC |
| // expected-error@+7 {{'f' requires 128 bit size '__float128' type support, but target 'nvptx64-unknown-unknown' does not support it}} |
| // expected-error@+6 {{expression requires 128 bit size '__float128' type support, but target 'nvptx64-unknown-unknown' does not support it}} |
| #else |
| // expected-error@+4 {{'f' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} |
| // expected-error@+3 {{expression requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} |
| #endif |
| T &operator+(T &b) { |
| f += b.a; |
| return *this; |
| } |
| }; |
| |
| struct T1 { |
| char a; |
| __int128 f; |
| __int128 f1; |
| char c; |
| T1() : a(12), f(15) {} |
| T1 &operator/(T1 &b) { |
| f /= b.a; |
| return *this; |
| } |
| }; |
| |
| #ifndef _ARCH_PPC |
| // expected-error@+2 {{'boo' requires 128 bit size '__float128' type support, but target 'nvptx64-unknown-unknown' does not support it}} |
| // expected-note@+1 2{{'boo' defined here}} |
| void boo(__float128 A) { return; } |
| #else |
| // expected-error@+2 {{'boo' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} |
| // expected-note@+1 2{{'boo' defined here}} |
| void boo(long double A) { return; } |
| #endif |
| #pragma omp declare target |
| T a = T(); |
| T f = a; |
| void foo(T a = T()) { |
| a = a + f; // expected-note {{called by 'foo'}} |
| #ifndef _ARCH_PPC |
| // expected-error@+5 {{'boo' requires 128 bit size '__float128' type support, but target 'nvptx64-unknown-unknown' does not support it}} |
| #else |
| // expected-error@+3 {{'boo' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} |
| #endif |
| // expected-note@+1 {{called by 'foo'}} |
| boo(0); |
| return; |
| } |
| T bar() { |
| return T(); |
| } |
| |
| void baz() { |
| T t = bar(); |
| } |
| T1 a1 = T1(); |
| T1 f1 = a1; |
| void foo1(T1 a = T1()) { |
| a = a / f1; |
| return; |
| } |
| T1 bar1() { |
| return T1(); |
| } |
| void baz1() { |
| T1 t = bar1(); |
| } |
| |
| inline void dead_inline_declare_target() { |
| long double *a, b = 0; |
| a = &b; |
| } |
| static void dead_static_declare_target() { |
| long double *a, b = 0; |
| a = &b; |
| } |
| template<bool> |
| void dead_template_declare_target() { |
| long double *a, b = 0; |
| a = &b; |
| } |
| |
| // expected-note@+2 {{'ld_return1a' defined here}} |
| // expected-error@+1 {{'ld_return1a' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} |
| long double ld_return1a() { return 0; } |
| // expected-note@+2 {{'ld_arg1a' defined here}} |
| // expected-error@+1 {{'ld_arg1a' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} |
| void ld_arg1a(long double ld) {} |
| |
| typedef long double ld_ty; |
| // expected-note@+2 {{'ld_return1b' defined here}} |
| // expected-error@+1 {{'ld_return1b' requires 128 bit size 'ld_ty' (aka 'long double') type support, but target 'nvptx64-unknown-unknown' does not support it}} |
| ld_ty ld_return1b() { return 0; } |
| // expected-note@+2 {{'ld_arg1b' defined here}} |
| // expected-error@+1 {{'ld_arg1b' requires 128 bit size 'ld_ty' (aka 'long double') type support, but target 'nvptx64-unknown-unknown' does not support it}} |
| void ld_arg1b(ld_ty ld) {} |
| |
| static long double ld_return1c() { return 0; } |
| static void ld_arg1c(long double ld) {} |
| |
| inline long double ld_return1d() { return 0; } |
| inline void ld_arg1d(long double ld) {} |
| |
| // expected-note@+1 {{'ld_return1e' defined here}} |
| static long double ld_return1e() { return 0; } |
| // expected-note@+1 {{'ld_arg1e' defined here}} |
| static void ld_arg1e(long double ld) {} |
| |
| // expected-note@+1 {{'ld_return1f' defined here}} |
| inline long double ld_return1f() { return 0; } |
| // expected-note@+1 {{'ld_arg1f' defined here}} |
| inline void ld_arg1f(long double ld) {} |
| |
| inline void ld_use1() { |
| long double ld = 0; |
| ld += 1; |
| } |
| static void ld_use2() { |
| long double ld = 0; |
| ld += 1; |
| } |
| |
| inline void ld_use3() { |
| // expected-note@+1 {{'ld' defined here}} |
| long double ld = 0; |
| // expected-error@+2 {{expression requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} |
| // expected-error@+1 {{'ld' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} |
| ld += 1; |
| } |
| static void ld_use4() { |
| // expected-note@+1 {{'ld' defined here}} |
| long double ld = 0; |
| // expected-error@+2 {{expression requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} |
| // expected-error@+1 {{'ld' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} |
| ld += 1; |
| } |
| |
| void external() { |
| // expected-error@+1 {{'ld_return1e' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} |
| void *p1 = reinterpret_cast<void*>(&ld_return1e); |
| // expected-error@+1 {{'ld_arg1e' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} |
| void *p2 = reinterpret_cast<void*>(&ld_arg1e); |
| // expected-error@+1 {{'ld_return1f' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} |
| void *p3 = reinterpret_cast<void*>(&ld_return1f); |
| // expected-error@+1 {{'ld_arg1f' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} |
| void *p4 = reinterpret_cast<void*>(&ld_arg1f); |
| // TODO: The error message "called by" is not great. |
| // expected-note@+1 {{called by 'external'}} |
| void *p5 = reinterpret_cast<void*>(&ld_use3); |
| // expected-note@+1 {{called by 'external'}} |
| void *p6 = reinterpret_cast<void*>(&ld_use4); |
| } |
| |
| #ifndef _ARCH_PPC |
| // expected-note@+2 {{'ld_return2a' defined here}} |
| // expected-error@+1 {{'ld_return2a' requires 128 bit size '__float128' type support, but target 'nvptx64-unknown-unknown' does not support it}} |
| __float128 ld_return2a() { return 0; } |
| // expected-note@+2 {{'ld_arg2a' defined here}} |
| // expected-error@+1 {{'ld_arg2a' requires 128 bit size '__float128' type support, but target 'nvptx64-unknown-unknown' does not support it}} |
| void ld_arg2a(__float128 ld) {} |
| |
| typedef __float128 fp128_ty; |
| // expected-note@+2 {{'ld_return2b' defined here}} |
| // expected-error@+1 {{'ld_return2b' requires 128 bit size 'fp128_ty' (aka '__float128') type support, but target 'nvptx64-unknown-unknown' does not support it}} |
| fp128_ty ld_return2b() { return 0; } |
| // expected-note@+2 {{'ld_arg2b' defined here}} |
| // expected-error@+1 {{'ld_arg2b' requires 128 bit size 'fp128_ty' (aka '__float128') type support, but target 'nvptx64-unknown-unknown' does not support it}} |
| void ld_arg2b(fp128_ty ld) {} |
| #endif |
| |
| #pragma omp end declare target |
| |
| // TODO: There should not be an error here, dead_inline is never emitted. |
| // expected-note@+1 {{'f' defined here}} |
| inline long double dead_inline(long double f) { |
| #pragma omp target map(f) |
| // expected-error@+1 {{'f' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} |
| f = 1; |
| return f; |
| } |
| |
| // TODO: There should not be an error here, dead_static is never emitted. |
| // expected-note@+1 {{'f' defined here}} |
| static long double dead_static(long double f) { |
| #pragma omp target map(f) |
| // expected-error@+1 {{'f' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} |
| f = 1; |
| return f; |
| } |
| |
| template<typename T> |
| long double dead_template(long double f) { |
| #pragma omp target map(f) |
| f = 1; |
| return f; |
| } |
| |
| #ifndef _ARCH_PPC |
| // expected-note@+1 {{'f' defined here}} |
| __float128 foo2(__float128 f) { |
| #pragma omp target map(f) |
| // expected-error@+1 {{'f' requires 128 bit size '__float128' type support, but target 'nvptx64-unknown-unknown' does not support it}} |
| f = 1; |
| return f; |
| } |
| #else |
| // expected-note@+1 {{'f' defined here}} |
| long double foo3(long double f) { |
| #pragma omp target map(f) |
| // expected-error@+1 {{'f' requires 128 bit size 'long double' type support, but target 'nvptx64-unknown-unknown' does not support it}} |
| f = 1; |
| return f; |
| } |
| #endif |
| |
| T foo3() { |
| T S; |
| #pragma omp target map(S) |
| S.a = 1; |
| return S; |
| } |
| |
| // Allow all sorts of stuff on host |
| #ifndef _ARCH_PPC |
| __float128 q, b; |
| __float128 c = q + b; |
| #else |
| long double q, b; |
| long double c = q + b; |
| #endif |
| |
| void hostFoo() { |
| boo(c - b); |
| } |
| |
| long double qa, qb; |
| decltype(qa + qb) qc; |
| double qd[sizeof(-(-(qc * 2)))]; |
| |
| struct A { }; |
| |
| template <bool> |
| struct A_type { typedef A type; }; |
| |
| template <class Sp, class Tp> |
| struct B { |
| enum { value = bool(Sp::value) || bool(Tp::value) }; |
| typedef typename A_type<value>::type type; |
| }; |
| |
| void bar(_ExtInt(66) a) { |
| auto b = a; |
| } |