clang/test/OpenMP/target_vla_messages.cpp

   1 // PowerPC supports VLAs.
   2 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-unknown-unknown -emit-llvm-bc %s -o %t-ppc-host-ppc.bc
   3 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host-ppc.bc -o %t-ppc-device.ll
   4
   5 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-unknown-unknown -emit-llvm-bc %s -o %t-ppc-host-ppc.bc
   6 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host-ppc.bc -o %t-ppc-device.ll
   7
   8 // Nvidia GPUs don't support VLAs.
   9 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host-nvptx.bc
  10 // RUN: %clang_cc1 -verify -DNO_VLA -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host-nvptx.bc -o %t-nvptx-device.ll
  11
  12 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host-nvptx.bc
  13
  14 #ifndef NO_VLA
  15 // expected-no-diagnostics
  16 #endif
  17
  18 #pragma omp declare target
  19 void declare(int arg) {
  20   int a[2];
  21 #ifdef NO_VLA
  22   // expected-error@+2 {{variable length arrays are not supported for the current target}}
  23 #endif
  24   int vla[arg];
  25 }
  26
  27 void declare_parallel_reduction(int arg) {
  28   int a[2];
  29
  30 #pragma omp parallel reduction(+: a)
  31   { }
  32
  33 #pragma omp parallel reduction(+: a[0:2])
  34   { }
  35
  36 #ifdef NO_VLA
  37   // expected-error@+3 {{cannot generate code for reduction on array section, which requires a variable length array}}
  38   // expected-note@+2 {{variable length arrays are not supported for the current target}}
  39 #endif
  40 #pragma omp parallel reduction(+: a[0:arg])
  41   { }
  42 }
  43 #pragma omp end declare target
  44
  45 template <typename T>
  46 void target_template(int arg) {
  47 #pragma omp target
  48   {
  49 #ifdef NO_VLA
  50     // expected-error@+2 {{variable length arrays are not supported for the current target}}
  51 #endif
  52     T vla[arg];
  53   }
  54 }
  55
  56 void target(int arg) {
  57 #pragma omp target
  58   {
  59 #ifdef NO_VLA
  60     // expected-error@+2 {{variable length arrays are not supported for the current target}}
  61 #endif
  62     int vla[arg];
  63   }
  64
  65 #pragma omp target
  66   {
  67 #pragma omp parallel
  68     {
  69 #ifdef NO_VLA
  70     // expected-error@+2 {{variable length arrays are not supported for the current target}}
  71 #endif
  72       int vla[arg];
  73     }
  74   }
  75
  76 #ifdef NO_VLA
  77     // expected-note@+2 {{in instantiation of function template specialization 'target_template<long>' requested here}}
  78 #endif
  79   target_template<long>(arg);
  80 }
  81
  82 void teams_reduction(int arg) {
  83   int a[2];
  84   int vla[arg];
  85
  86 #pragma omp target map(a)
  87 #pragma omp teams reduction(+: a)
  88   { }
  89
  90 #ifdef NO_VLA
  91   // expected-error@+4 {{cannot generate code for reduction on variable length array}}
  92   // expected-note@+3 {{variable length arrays are not supported for the current target}}
  93 #endif
  94 #pragma omp target map(vla)
  95 #pragma omp teams reduction(+: vla)
  96   { }
  97
  98 #pragma omp target map(a[0:2])
  99 #pragma omp teams reduction(+: a[0:2])
 100   { }
 101
 102 #pragma omp target map(vla[0:2])
 103 #pragma omp teams reduction(+: vla[0:2])
 104   { }
 105
 106 #ifdef NO_VLA
 107   // expected-error@+4 {{cannot generate code for reduction on array section, which requires a variable length array}}
 108   // expected-note@+3 {{variable length arrays are not supported for the current target}}
 109 #endif
 110 #pragma omp target map(a[0:arg])
 111 #pragma omp teams reduction(+: a[0:arg])
 112   { }
 113
 114 #ifdef NO_VLA
 115   // expected-error@+4 {{cannot generate code for reduction on array section, which requires a variable length array}}
 116   // expected-note@+3 {{variable length arrays are not supported for the current target}}
 117 #endif
 118 #pragma omp target map(vla[0:arg])
 119 #pragma omp teams reduction(+: vla[0:arg])
 120   { }
 121 }
 122
 123 void parallel_reduction(int arg) {
 124   int a[2];
 125   int vla[arg];
 126
 127 #pragma omp target map(a)
 128 #pragma omp parallel reduction(+: a)
 129   { }
 130
 131 #ifdef NO_VLA
 132   // expected-error@+4 {{cannot generate code for reduction on variable length array}}
 133   // expected-note@+3 {{variable length arrays are not supported for the current target}}
 134 #endif
 135 #pragma omp target map(vla)
 136 #pragma omp parallel reduction(+: vla)
 137   { }
 138
 139 #pragma omp target map(a[0:2])
 140 #pragma omp parallel reduction(+: a[0:2])
 141   { }
 142
 143 #pragma omp target map(vla[0:2])
 144 #pragma omp parallel reduction(+: vla[0:2])
 145   { }
 146
 147 #ifdef NO_VLA
 148   // expected-error@+4 {{cannot generate code for reduction on array section, which requires a variable length array}}
 149   // expected-note@+3 {{variable length arrays are not supported for the current target}}
 150 #endif
 151 #pragma omp target map(a[0:arg])
 152 #pragma omp parallel reduction(+: a[0:arg])
 153   { }
 154
 155 #ifdef NO_VLA
 156   // expected-error@+4 {{cannot generate code for reduction on array section, which requires a variable length array}}
 157   // expected-note@+3 {{variable length arrays are not supported for the current target}}
 158 #endif
 159 #pragma omp target map(vla[0:arg])
 160 #pragma omp parallel reduction(+: vla[0:arg])
 161   { }
 162 }
 163
 164 void for_reduction(int arg) {
 165   int a[2];
 166   int vla[arg];
 167
 168 #pragma omp target map(a)
 169 #pragma omp parallel
 170 #pragma omp for reduction(+: a)
 171   for (int i = 0; i < arg; i++) ;
 172
 173 #ifdef NO_VLA
 174   // expected-error@+5 {{cannot generate code for reduction on variable length array}}
 175   // expected-note@+4 {{variable length arrays are not supported for the current target}}
 176 #endif
 177 #pragma omp target map(vla)
 178 #pragma omp parallel
 179 #pragma omp for reduction(+: vla)
 180   for (int i = 0; i < arg; i++) ;
 181
 182 #pragma omp target map(a[0:2])
 183 #pragma omp parallel
 184 #pragma omp for reduction(+: a[0:2])
 185   for (int i = 0; i < arg; i++) ;
 186
 187 #pragma omp target map(vla[0:2])
 188 #pragma omp parallel
 189 #pragma omp for reduction(+: vla[0:2])
 190   for (int i = 0; i < arg; i++) ;
 191
 192 #ifdef NO_VLA
 193   // expected-error@+5 {{cannot generate code for reduction on array section, which requires a variable length array}}
 194   // expected-note@+4 {{variable length arrays are not supported for the current target}}
 195 #endif
 196 #pragma omp target map(a[0:arg])
 197 #pragma omp parallel
 198 #pragma omp for reduction(+: a[0:arg])
 199   for (int i = 0; i < arg; i++) ;
 200
 201 #ifdef NO_VLA
 202   // expected-error@+5 {{cannot generate code for reduction on array section, which requires a variable length array}}
 203   // expected-note@+4 {{variable length arrays are not supported for the current target}}
 204 #endif
 205 #pragma omp target map(vla[0:arg])
 206 #pragma omp parallel
 207 #pragma omp for reduction(+: vla[0:arg])
 208   for (int i = 0; i < arg; i++) ;
 209 #ifdef NO_VLA
 210   // expected-error@+3 {{cannot generate code for reduction on array section, which requires a variable length array}}
 211   // expected-note@+2 {{variable length arrays are not supported for the current target}}
 212 #endif
 213 #pragma omp target reduction(+ : vla[0:arg])
 214   for (int i = 0; i < arg; i++) ;
 215 }