summaryrefslogtreecommitdiff
path: root/test/OpenMP/target_vla_messages.cpp
diff options
context:
space:
mode:
Diffstat (limited to 'test/OpenMP/target_vla_messages.cpp')
-rw-r--r--test/OpenMP/target_vla_messages.cpp201
1 files changed, 201 insertions, 0 deletions
diff --git a/test/OpenMP/target_vla_messages.cpp b/test/OpenMP/target_vla_messages.cpp
new file mode 100644
index 000000000000..c970d0def661
--- /dev/null
+++ b/test/OpenMP/target_vla_messages.cpp
@@ -0,0 +1,201 @@
+// PowerPC supports VLAs.
+// 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
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host-ppc.bc -o %t-ppc-device.ll
+
+// Nvidia GPUs don't support VLAs.
+// 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
+// RUN: %clang_cc1 -verify -DNO_VLA -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host-nvptx.bc -o %t-nvptx-device.ll
+
+#ifndef NO_VLA
+// expected-no-diagnostics
+#endif
+
+#pragma omp declare target
+void declare(int arg) {
+ int a[2];
+#ifdef NO_VLA
+ // expected-error@+2 {{variable length arrays are not supported for the current target}}
+#endif
+ int vla[arg];
+}
+
+void declare_parallel_reduction(int arg) {
+ int a[2];
+
+#pragma omp parallel reduction(+: a)
+ { }
+
+#pragma omp parallel reduction(+: a[0:2])
+ { }
+
+#ifdef NO_VLA
+ // expected-error@+3 {{cannot generate code for reduction on array section, which requires a variable length array}}
+ // expected-note@+2 {{variable length arrays are not supported for the current target}}
+#endif
+#pragma omp parallel reduction(+: a[0:arg])
+ { }
+}
+#pragma omp end declare target
+
+template <typename T>
+void target_template(int arg) {
+#pragma omp target
+ {
+#ifdef NO_VLA
+ // expected-error@+2 {{variable length arrays are not supported for the current target}}
+#endif
+ T vla[arg];
+ }
+}
+
+void target(int arg) {
+#pragma omp target
+ {
+#ifdef NO_VLA
+ // expected-error@+2 {{variable length arrays are not supported for the current target}}
+#endif
+ int vla[arg];
+ }
+
+#pragma omp target
+ {
+#pragma omp parallel
+ {
+#ifdef NO_VLA
+ // expected-error@+2 {{variable length arrays are not supported for the current target}}
+#endif
+ int vla[arg];
+ }
+ }
+
+ target_template<long>(arg);
+}
+
+void teams_reduction(int arg) {
+ int a[2];
+ int vla[arg];
+
+#pragma omp target map(a)
+#pragma omp teams reduction(+: a)
+ { }
+
+#ifdef NO_VLA
+ // expected-error@+4 {{cannot generate code for reduction on variable length array}}
+ // expected-note@+3 {{variable length arrays are not supported for the current target}}
+#endif
+#pragma omp target map(vla)
+#pragma omp teams reduction(+: vla)
+ { }
+
+#pragma omp target map(a[0:2])
+#pragma omp teams reduction(+: a[0:2])
+ { }
+
+#pragma omp target map(vla[0:2])
+#pragma omp teams reduction(+: vla[0:2])
+ { }
+
+#ifdef NO_VLA
+ // expected-error@+4 {{cannot generate code for reduction on array section, which requires a variable length array}}
+ // expected-note@+3 {{variable length arrays are not supported for the current target}}
+#endif
+#pragma omp target map(a[0:arg])
+#pragma omp teams reduction(+: a[0:arg])
+ { }
+
+#ifdef NO_VLA
+ // expected-error@+4 {{cannot generate code for reduction on array section, which requires a variable length array}}
+ // expected-note@+3 {{variable length arrays are not supported for the current target}}
+#endif
+#pragma omp target map(vla[0:arg])
+#pragma omp teams reduction(+: vla[0:arg])
+ { }
+}
+
+void parallel_reduction(int arg) {
+ int a[2];
+ int vla[arg];
+
+#pragma omp target map(a)
+#pragma omp parallel reduction(+: a)
+ { }
+
+#ifdef NO_VLA
+ // expected-error@+4 {{cannot generate code for reduction on variable length array}}
+ // expected-note@+3 {{variable length arrays are not supported for the current target}}
+#endif
+#pragma omp target map(vla)
+#pragma omp parallel reduction(+: vla)
+ { }
+
+#pragma omp target map(a[0:2])
+#pragma omp parallel reduction(+: a[0:2])
+ { }
+
+#pragma omp target map(vla[0:2])
+#pragma omp parallel reduction(+: vla[0:2])
+ { }
+
+#ifdef NO_VLA
+ // expected-error@+4 {{cannot generate code for reduction on array section, which requires a variable length array}}
+ // expected-note@+3 {{variable length arrays are not supported for the current target}}
+#endif
+#pragma omp target map(a[0:arg])
+#pragma omp parallel reduction(+: a[0:arg])
+ { }
+
+#ifdef NO_VLA
+ // expected-error@+4 {{cannot generate code for reduction on array section, which requires a variable length array}}
+ // expected-note@+3 {{variable length arrays are not supported for the current target}}
+#endif
+#pragma omp target map(vla[0:arg])
+#pragma omp parallel reduction(+: vla[0:arg])
+ { }
+}
+
+void for_reduction(int arg) {
+ int a[2];
+ int vla[arg];
+
+#pragma omp target map(a)
+#pragma omp parallel
+#pragma omp for reduction(+: a)
+ for (int i = 0; i < arg; i++) ;
+
+#ifdef NO_VLA
+ // expected-error@+5 {{cannot generate code for reduction on variable length array}}
+ // expected-note@+4 {{variable length arrays are not supported for the current target}}
+#endif
+#pragma omp target map(vla)
+#pragma omp parallel
+#pragma omp for reduction(+: vla)
+ for (int i = 0; i < arg; i++) ;
+
+#pragma omp target map(a[0:2])
+#pragma omp parallel
+#pragma omp for reduction(+: a[0:2])
+ for (int i = 0; i < arg; i++) ;
+
+#pragma omp target map(vla[0:2])
+#pragma omp parallel
+#pragma omp for reduction(+: vla[0:2])
+ for (int i = 0; i < arg; i++) ;
+
+#ifdef NO_VLA
+ // expected-error@+5 {{cannot generate code for reduction on array section, which requires a variable length array}}
+ // expected-note@+4 {{variable length arrays are not supported for the current target}}
+#endif
+#pragma omp target map(a[0:arg])
+#pragma omp parallel
+#pragma omp for reduction(+: a[0:arg])
+ for (int i = 0; i < arg; i++) ;
+
+#ifdef NO_VLA
+ // expected-error@+5 {{cannot generate code for reduction on array section, which requires a variable length array}}
+ // expected-note@+4 {{variable length arrays are not supported for the current target}}
+#endif
+#pragma omp target map(vla[0:arg])
+#pragma omp parallel
+#pragma omp for reduction(+: vla[0:arg])
+ for (int i = 0; i < arg; i++) ;
+}