151 lines · c
1// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc2// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized3// RUN: %clang_cc1 -verify -DDIAGS -DIMMEDIATE -fopenmp -fopenmp-version=45 -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized4// RUN: %clang_cc1 -verify -DDIAGS -DDELAYED -fopenmp -fopenmp-version=45 -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized5// RUN: %clang_cc1 -fopenmp -x c -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc6// RUN: %clang_cc1 -verify=expected,omp5 -fopenmp -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized7// RUN: %clang_cc1 -verify=expected,omp5 -DDIAGS -DOMP5 -DIMMEDIATE -fopenmp -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized8// RUN: %clang_cc1 -verify=expected,omp5 -DDIAGS -DOMP5 -DDELAYED -fopenmp -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized9// REQUIRES: x86-registered-target10// REQUIRES: nvptx-registered-target11 12#ifndef DIAGS13// expected-no-diagnostics14#endif // DIAGS15 16#ifdef OMP517void bar(int r) {18#ifdef IMMEDIATE19// omp5-error@+4 {{invalid input constraint 'mx' in asm}}20#endif // IMMEDIATE21 __asm__("PR3908 %[lf] %[xx] %[li] %[r]"22 : [ r ] "+r"(r)23 : [ lf ] "mx"(0), [ li ] "mr"(0), [ xx ] "x"((double)(0)));24}25#ifdef IMMEDIATE26#pragma omp declare target to(bar) device_type(nohost)27#else28#pragma omp declare target to(bar) device_type(host)29#endif // IMMEDIATE30#endif // OMP531 32void foo(int r) {33#ifdef IMMEDIATE34// expected-error@+4 {{invalid input constraint 'mx' in asm}}35#endif // IMMEDIATE36 __asm__("PR3908 %[lf] %[xx] %[li] %[r]"37 : [ r ] "+r"(r)38 : [ lf ] "mx"(0), [ li ] "mr"(0), [ xx ] "x"((double)(0)));39}40#ifdef IMMEDIATE41#pragma omp declare target to(foo)42#endif //IMMEDIATE43 44#ifdef IMMEDIATE45#pragma omp declare target46#endif //IMMEDIATE47void t1(int r) {48#ifdef DIAGS49// expected-error@+4 {{invalid input constraint 'mx' in asm}}50#endif // DIAGS51 __asm__("PR3908 %[lf] %[xx] %[li] %[r]"52 : [ r ] "+r"(r)53 : [ lf ] "mx"(0), [ li ] "mr"(0), [ xx ] "x"((double)(0)));54}55 56unsigned t2(signed char input) {57 unsigned output;58#ifdef DIAGS59// expected-error@+3 {{invalid output constraint '=a' in asm}}60#endif // DIAGS61 __asm__("xyz"62 : "=a"(output)63 : "0"(input));64 return output;65}66 67double t3(double x) {68 register long double result;69#ifdef DIAGS70// expected-error@+3 {{invalid output constraint '=t' in asm}}71#endif // DIAGS72 __asm __volatile("frndint"73 : "=t"(result)74 : "0"(x));75 return result;76}77 78unsigned char t4(unsigned char a, unsigned char b) {79 unsigned int la = a;80 unsigned int lb = b;81 unsigned int bigres;82 unsigned char res;83#ifdef DIAGS84// expected-error@+3 {{invalid output constraint '=la' in asm}}85#endif // DIAGS86 __asm__("0:\n1:\n"87 : [ bigres ] "=la"(bigres)88 : [ la ] "0"(la), [ lb ] "c"(lb)89 : "edx", "cc");90 res = bigres;91 return res;92}93 94void t5(void) {95#ifdef DIAGS96// expected-error@+6 {{unknown register name 'st' in asm}}97#endif // DIAGS98 __asm__ __volatile__(99 "finit"100 :101 :102 : "st", "st(1)", "st(2)", "st(3)",103 "st(4)", "st(5)", "st(6)", "st(7)",104 "fpsr", "fpcr");105}106 107typedef long long __m256i __attribute__((__vector_size__(32)));108void t6(__m256i *p) {109#ifdef DIAGS110// expected-error@+3 {{unknown register name 'ymm0' in asm}}111#endif // DIAGS112 __asm__ volatile("vmovaps %0, %%ymm0" ::"m"(*(__m256i *)p)113 : "ymm0");114}115#ifdef IMMEDIATE116#pragma omp end declare target117#endif //IMMEDIATE118 119int main() {120#ifdef DELAYED121#pragma omp target122#endif // DELAYED123 {124#ifdef DELAYED125// expected-note@+2 {{called by 'main'}}126#endif // DELAYED127 t1(0);128#ifdef DELAYED129// expected-note@+2 {{called by 'main'}}130#endif // DELAYED131 t2(0);132#ifdef DELAYED133// expected-note@+2 {{called by 'main'}}134#endif // DELAYED135 t3(0);136#ifdef DELAYED137// expected-note@+2 {{called by 'main'}}138#endif // DELAYED139 t4(0, 0);140#ifdef DELAYED141// expected-note@+2 {{called by 'main'}}142#endif // DELAYED143 t5();144#ifdef DELAYED145// expected-note@+2 {{called by 'main'}}146#endif // DELAYED147 t6(0);148 }149 return 0;150}151