clang/test/OpenMP/nvptx_asm_delayed_diags.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.bc
   2 // 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 -Wuninitialized
   3 // 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 -Wuninitialized
   4 // 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 -Wuninitialized
   5 // RUN: %clang_cc1 -fopenmp -x c -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
   6 // 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 -Wuninitialized
   7 // 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 -Wuninitialized
   8 // 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 -Wuninitialized
   9 // REQUIRES: x86-registered-target
  10 // REQUIRES: nvptx-registered-target
  11
  12 #ifndef DIAGS
  13 // expected-no-diagnostics
  14 #endif // DIAGS
  15
  16 #ifdef OMP5
  17 void bar(int r) {
  18 #ifdef IMMEDIATE
  19 // omp5-error@+4 {{invalid input constraint 'mx' in asm}}
  20 #endif // IMMEDIATE
  21   __asm__("PR3908 %[lf] %[xx] %[li] %[r]"
  22           : [ r ] "+r"(r)
  23           : [ lf ] "mx"(0), [ li ] "mr"(0), [ xx ] "x"((double)(0)));
  24 }
  25 #ifdef IMMEDIATE
  26 #pragma omp declare target to(bar) device_type(nohost)
  27 #else
  28 #pragma omp declare target to(bar) device_type(host)
  29 #endif // IMMEDIATE
  30 #endif // OMP5
  31
  32 void foo(int r) {
  33 #ifdef IMMEDIATE
  34 // expected-error@+4 {{invalid input constraint 'mx' in asm}}
  35 #endif // IMMEDIATE
  36   __asm__("PR3908 %[lf] %[xx] %[li] %[r]"
  37           : [ r ] "+r"(r)
  38           : [ lf ] "mx"(0), [ li ] "mr"(0), [ xx ] "x"((double)(0)));
  39 }
  40 #ifdef IMMEDIATE
  41 #pragma omp declare target to(foo)
  42 #endif //IMMEDIATE
  43
  44 #ifdef IMMEDIATE
  45 #pragma omp declare target
  46 #endif //IMMEDIATE
  47 void t1(int r) {
  48 #ifdef DIAGS
  49 // expected-error@+4 {{invalid input constraint 'mx' in asm}}
  50 #endif // DIAGS
  51   __asm__("PR3908 %[lf] %[xx] %[li] %[r]"
  52           : [ r ] "+r"(r)
  53           : [ lf ] "mx"(0), [ li ] "mr"(0), [ xx ] "x"((double)(0)));
  54 }
  55
  56 unsigned t2(signed char input) {
  57   unsigned output;
  58 #ifdef DIAGS
  59 // expected-error@+3 {{invalid output constraint '=a' in asm}}
  60 #endif // DIAGS
  61   __asm__("xyz"
  62           : "=a"(output)
  63           : "0"(input));
  64   return output;
  65 }
  66
  67 double t3(double x) {
  68   register long double result;
  69 #ifdef DIAGS
  70 // expected-error@+3 {{invalid output constraint '=t' in asm}}
  71 #endif // DIAGS
  72   __asm __volatile("frndint"
  73                    : "=t"(result)
  74                    : "0"(x));
  75   return result;
  76 }
  77
  78 unsigned 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 DIAGS
  84 // expected-error@+3 {{invalid output constraint '=la' in asm}}
  85 #endif // DIAGS
  86   __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
  94 void t5(void) {
  95 #ifdef DIAGS
  96 // expected-error@+6 {{unknown register name 'st' in asm}}
  97 #endif // DIAGS
  98   __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
 107 typedef long long __m256i __attribute__((__vector_size__(32)));
 108 void t6(__m256i *p) {
 109 #ifdef DIAGS
 110 // expected-error@+3 {{unknown register name 'ymm0' in asm}}
 111 #endif // DIAGS
 112   __asm__ volatile("vmovaps  %0, %%ymm0" ::"m"(*(__m256i *)p)
 113                    : "ymm0");
 114 }
 115 #ifdef IMMEDIATE
 116 #pragma omp end declare target
 117 #endif //IMMEDIATE
 118
 119 int main() {
 120 #ifdef DELAYED
 121 #pragma omp target
 122 #endif // DELAYED
 123   {
 124 #ifdef DELAYED
 125 // expected-note@+2 {{called by 'main'}}
 126 #endif // DELAYED
 127     t1(0);
 128 #ifdef DELAYED
 129 // expected-note@+2 {{called by 'main'}}
 130 #endif // DELAYED
 131     t2(0);
 132 #ifdef DELAYED
 133 // expected-note@+2 {{called by 'main'}}
 134 #endif // DELAYED
 135     t3(0);
 136 #ifdef DELAYED
 137 // expected-note@+2 {{called by 'main'}}
 138 #endif // DELAYED
 139     t4(0, 0);
 140 #ifdef DELAYED
 141 // expected-note@+2 {{called by 'main'}}
 142 #endif // DELAYED
 143     t5();
 144 #ifdef DELAYED
 145 // expected-note@+2 {{called by 'main'}}
 146 #endif // DELAYED
 147     t6(0);
 148   }
 149   return 0;
 150 }