aboutsummaryrefslogtreecommitdiff
path: root/test/OpenMP/nvptx_target_cuda_mode_messages.cpp
diff options
context:
space:
mode:
authorDimitry Andric <dim@FreeBSD.org>2018-07-28 11:06:01 +0000
committerDimitry Andric <dim@FreeBSD.org>2018-07-28 11:06:01 +0000
commit486754660bb926339aefcf012a3f848592babb8b (patch)
treeecdbc446c9876f4f120f701c243373cd3cb43db3 /test/OpenMP/nvptx_target_cuda_mode_messages.cpp
parent55e6d896ad333f07bb3b1ba487df214fc268a4ab (diff)
downloadsrc-vendor/clang/clang-trunk-r338150.tar.gz
src-vendor/clang/clang-trunk-r338150.zip
Vendor import of clang trunk r338150:vendor/clang/clang-trunk-r338150
Diffstat (limited to 'test/OpenMP/nvptx_target_cuda_mode_messages.cpp')
-rw-r--r--test/OpenMP/nvptx_target_cuda_mode_messages.cpp108
1 files changed, 108 insertions, 0 deletions
diff --git a/test/OpenMP/nvptx_target_cuda_mode_messages.cpp b/test/OpenMP/nvptx_target_cuda_mode_messages.cpp
new file mode 100644
index 000000000000..eecc26c6181a
--- /dev/null
+++ b/test/OpenMP/nvptx_target_cuda_mode_messages.cpp
@@ -0,0 +1,108 @@
+// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-cuda-mode -fopenmp-host-ir-file-path %t-ppc-host.bc -o -
+
+template <typename tx, typename ty>
+struct TT {
+ tx X;
+ ty Y;
+};
+
+int foo(int n, double *ptr) {
+ int a = 0;
+ short aa = 0;
+ float b[10];
+ double c[5][10];
+ TT<long long, char> d;
+
+#pragma omp target firstprivate(a) map(tofrom: b) // expected-note 2 {{defined as threadprivate or thread local}}
+ {
+ int c; // expected-note {{defined as threadprivate or thread local}}
+#pragma omp parallel shared(a, b, c, aa) // expected-error 3 {{threadprivate or thread local variable cannot be shared}}
+ b[a] = a;
+#pragma omp parallel for
+ for (int i = 0; i < 10; ++i) // expected-note {{defined as threadprivate or thread local}}
+#pragma omp parallel shared(i) // expected-error {{threadprivate or thread local variable cannot be shared}}
+ ++i;
+ }
+
+#pragma omp target map(aa, b, c, d)
+ {
+ int e; // expected-note {{defined as threadprivate or thread local}}
+#pragma omp parallel private(b, e) // expected-error {{threadprivate or thread local variable cannot be private}}
+ {
+ aa += 1;
+ b[2] = 1.0;
+ c[1][2] = 1.0;
+ d.X = 1;
+ d.Y = 1;
+ }
+ }
+
+#pragma omp target private(ptr)
+ {
+ ptr[0]++;
+ }
+
+ return a;
+}
+
+template <typename tx>
+tx ftemplate(int n) {
+ tx a = 0;
+ tx b[10];
+
+#pragma omp target reduction(+ \
+ : a, b) // expected-note {{defined as threadprivate or thread local}}
+ {
+ int e; // expected-note {{defined as threadprivate or thread local}}
+#pragma omp parallel shared(a, e) // expected-error 2 {{threadprivate or thread local variable cannot be shared}}
+ a += 1;
+ b[2] += 1;
+ }
+
+ return a;
+}
+
+static int fstatic(int n) {
+ int a = 0;
+ char aaa = 0;
+ int b[10];
+
+#pragma omp target firstprivate(a, aaa, b)
+ {
+ a += 1;
+ aaa += 1;
+ b[2] += 1;
+ }
+
+ return a;
+}
+
+struct S1 {
+ double a;
+
+ int r1(int n) {
+ int b = n + 1;
+
+#pragma omp target firstprivate(b) // expected-note {{defined as threadprivate or thread local}}
+ {
+ int c; // expected-note {{defined as threadprivate or thread local}}
+#pragma omp parallel shared(b, c) // expected-error 2 {{threadprivate or thread local variable cannot be shared}}
+ this->a = (double)b + 1.5;
+ }
+
+ return (int)b;
+ }
+};
+
+int bar(int n, double *ptr) {
+ int a = 0;
+ a += foo(n, ptr);
+ S1 S;
+ a += S.r1(n);
+ a += fstatic(n);
+ a += ftemplate<int>(n); // expected-note {{in instantiation of function template specialization 'ftemplate<int>' requested here}}
+
+ return a;
+}
+