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 | |
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> |
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 | |
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 | |
84 | 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 | |
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 | |