1 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
2 // 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 -
3
4 template <typename tx, typename ty>
5 struct TT {
6 tx X;
7 ty Y;
8 };
9
foo(int n,double * ptr)10 int foo(int n, double *ptr) {
11 int a = 0;
12 short aa = 0;
13 float b[10];
14 double c[5][10];
15 TT<long long, char> d;
16
17 #pragma omp target firstprivate(a) map(tofrom: b) // expected-note 2 {{defined as threadprivate or thread local}}
18 {
19 int c; // expected-note {{defined as threadprivate or thread local}}
20 #pragma omp parallel shared(a, b, c, aa) // expected-error 3 {{threadprivate or thread local variable cannot be shared}}
21 b[a] = a;
22 #pragma omp parallel for
23 for (int i = 0; i < 10; ++i) // expected-note {{defined as threadprivate or thread local}}
24 #pragma omp parallel shared(i) // expected-error {{threadprivate or thread local variable cannot be shared}}
25 ++i;
26 }
27
28 #pragma omp target map(aa, b, c, d)
29 {
30 int e; // expected-note {{defined as threadprivate or thread local}}
31 #pragma omp parallel private(b, e) // expected-error {{threadprivate or thread local variable cannot be private}}
32 {
33 aa += 1;
34 b[2] = 1.0;
35 c[1][2] = 1.0;
36 d.X = 1;
37 d.Y = 1;
38 }
39 }
40
41 #pragma omp target private(ptr)
42 {
43 ptr[0]++;
44 }
45
46 return a;
47 }
48
49 template <typename tx>
ftemplate(int n)50 tx ftemplate(int n) {
51 tx a = 0;
52 tx b[10];
53
54 #pragma omp target reduction(+ \
55 : a, b) // expected-note {{defined as threadprivate or thread local}}
56 {
57 int e; // expected-note {{defined as threadprivate or thread local}}
58 #pragma omp parallel shared(a, e) // expected-error 2 {{threadprivate or thread local variable cannot be shared}}
59 a += 1;
60 b[2] += 1;
61 }
62
63 return a;
64 }
65
fstatic(int n)66 static int fstatic(int n) {
67 int a = 0;
68 char aaa = 0;
69 int b[10];
70
71 #pragma omp target firstprivate(a, aaa, b)
72 {
73 a += 1;
74 aaa += 1;
75 b[2] += 1;
76 }
77
78 return a;
79 }
80
81 struct S1 {
82 double a;
83
r1S184 int r1(int n) {
85 int b = n + 1;
86
87 #pragma omp target firstprivate(b) // expected-note {{defined as threadprivate or thread local}}
88 {
89 int c; // expected-note {{defined as threadprivate or thread local}}
90 #pragma omp parallel shared(b, c) // expected-error 2 {{threadprivate or thread local variable cannot be shared}}
91 this->a = (double)b + 1.5;
92 }
93
94 return (int)b;
95 }
96 };
97
bar(int n,double * ptr)98 int bar(int n, double *ptr) {
99 int a = 0;
100 a += foo(n, ptr);
101 S1 S;
102 a += S.r1(n);
103 a += fstatic(n);
104 a += ftemplate<int>(n); // expected-note {{in instantiation of function template specialization 'ftemplate<int>' requested here}}
105
106 return a;
107 }
108
109