1 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -ast-print %s -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP45
2 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -emit-pch -o %t %s -Wno-openmp-mapping
3 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -std=c++11 -include-pch %t -verify %s -ast-print -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP45
4 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -DOMP5 -ast-print %s -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP50
5 // RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -DOMP5 -x c++ -std=c++11 -emit-pch -o %t %s -Wno-openmp-mapping
6 // RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -DOMP5 -std=c++11 -include-pch %t -verify %s -ast-print -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP50
7 // RUN: %clang_cc1 -verify -fopenmp -DOMP51 -ast-print %s -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP51
8 // RUN: %clang_cc1 -fopenmp -DOMP51 -x c++ -std=c++11 -emit-pch -o %t %s -Wno-openmp-mapping
9 // RUN: %clang_cc1 -fopenmp -DOMP51 -std=c++11 -include-pch %t -verify %s -ast-print -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP51
10
11 // RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -ast-print %s -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP45
12 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -emit-pch -o %t %s -Wno-openmp-mapping
13 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -std=c++11 -include-pch %t -verify %s -ast-print -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP45
14 // RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=50 -DOMP5 -ast-print %s -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP50
15 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -DOMP5 -x c++ -std=c++11 -emit-pch -o %t %s -Wno-openmp-mapping
16 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -DOMP5 -std=c++11 -include-pch %t -verify %s -ast-print -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP50
17 // RUN: %clang_cc1 -verify -fopenmp-simd -DOMP51 -ast-print %s -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP51
18 // RUN: %clang_cc1 -fopenmp-simd -DOMP51 -x c++ -std=c++11 -emit-pch -o %t %s -Wno-openmp-mapping
19 // RUN: %clang_cc1 -fopenmp-simd -DOMP51 -std=c++11 -include-pch %t -verify %s -ast-print -Wno-openmp-mapping | FileCheck %s --check-prefix=CHECK --check-prefix=OMP51
20 // expected-no-diagnostics
21
22 #ifndef HEADER
23 #define HEADER
24
25 typedef void **omp_allocator_handle_t;
26 extern const omp_allocator_handle_t omp_null_allocator;
27 extern const omp_allocator_handle_t omp_default_mem_alloc;
28 extern const omp_allocator_handle_t omp_large_cap_mem_alloc;
29 extern const omp_allocator_handle_t omp_const_mem_alloc;
30 extern const omp_allocator_handle_t omp_high_bw_mem_alloc;
31 extern const omp_allocator_handle_t omp_low_lat_mem_alloc;
32 extern const omp_allocator_handle_t omp_cgroup_mem_alloc;
33 extern const omp_allocator_handle_t omp_pteam_mem_alloc;
34 extern const omp_allocator_handle_t omp_thread_mem_alloc;
35
foo()36 void foo() {}
37
38 struct S {
SS39 S(): a(0) {}
SS40 S(int v) : a(v) {}
41 int a;
42 typedef int type;
43 };
44
45 template <typename T>
46 class S7 : public T {
47 protected:
48 T a;
S7()49 S7() : a(0) {}
50
51 public:
S7(typename T::type v)52 S7(typename T::type v) : a(v) {
53 #pragma omp target teams distribute simd private(a) private(this->a) private(T::a)
54 for (int k = 0; k < a.a; ++k)
55 ++this->a.a;
56 }
operator =(S7 & s)57 S7 &operator=(S7 &s) {
58 int k;
59 #pragma omp target teams distribute simd allocate(a) private(a) private(this->a) linear(k) allocate(k)
60 for (k = 0; k < s.a.a; ++k)
61 ++s.a.a;
62 return *this;
63 }
foo()64 void foo() {
65 int b, argv, d, c, e, f;
66 #pragma omp target teams distribute simd private(b), firstprivate(argv) shared(d) reduction(+:c) reduction(max:e) num_teams(f) thread_limit(d)
67 for (int k = 0; k < a.a; ++k)
68 ++a.a;
69 }
bar()70 void bar() {
71 int arr[10];
72 const int alen = 16;
73 const int slen1 = 8;
74 const int slen2 = 8;
75 #pragma omp target teams distribute simd simdlen(slen1) safelen(slen2) aligned(arr:alen)
76 for (int k = 0; k < a.a; ++k)
77 ++a.a;
78 }
79 };
80 // CHECK: #pragma omp target teams distribute simd private(this->a) private(this->a) private(T::a)
81 // CHECK: #pragma omp target teams distribute simd allocate(this->a) private(this->a) private(this->a) linear(k) allocate(k)
82 // CHECK: #pragma omp target teams distribute simd private(b) firstprivate(argv) shared(d) reduction(+: c) reduction(max: e) num_teams(f) thread_limit(d)
83 // CHECK: #pragma omp target teams distribute simd simdlen(slen1) safelen(slen2) aligned(arr: alen)
84 // CHECK: #pragma omp target teams distribute simd private(this->a) private(this->a) private(this->S::a)
85
86 class S8 : public S7<S> {
S8()87 S8() {}
88
89 public:
S8(int v)90 S8(int v) : S7<S>(v){
91 #pragma omp target teams distribute simd private(a) private(this->a) private(S7<S>::a)
92 for (int k = 0; k < a.a; ++k)
93 ++this->a.a;
94 }
operator =(S8 & s)95 S8 &operator=(S8 &s) {
96 #pragma omp target teams distribute simd private(a) private(this->a)
97 for (int k = 0; k < s.a.a; ++k)
98 ++s.a.a;
99 return *this;
100 }
bar()101 void bar() {
102 int b, argv, d, c, e, f;
103 #pragma omp target teams distribute simd private(b), firstprivate(argv) shared(d) reduction(+:c) reduction(max:e) num_teams(f) thread_limit(d)
104 for (int k = 0; k < a.a; ++k)
105 ++a.a;
106 }
foo()107 void foo() {
108 const int alen = 16;
109 const int slen1 = 8;
110 const int slen2 = 8;
111 int arr[10];
112 #pragma omp target teams distribute simd simdlen(slen1) safelen(slen2) aligned(arr:alen)
113 for (int k = 0; k < a.a; ++k)
114 ++a.a;
115 }
116 };
117 // CHECK: #pragma omp target teams distribute simd private(this->a) private(this->a) private(this->S7<S>::a)
118 // CHECK: #pragma omp target teams distribute simd private(this->a) private(this->a)
119 // CHECK: #pragma omp target teams distribute simd private(b) firstprivate(argv) shared(d) reduction(+: c) reduction(max: e) num_teams(f) thread_limit(d)
120 // CHECK: #pragma omp target teams distribute simd simdlen(slen1) safelen(slen2) aligned(arr: alen)
121
122 template <class T, int N>
tmain(T argc)123 T tmain(T argc) {
124 T b = argc, c, d, e, f, g;
125 static T a;
126 // CHECK: static T a;
127 const T clen = 5;
128 const T alen = 16;
129 int arr[10];
130 #pragma omp target teams distribute simd
131 for (int i=0; i < 2; ++i)
132 a = 2;
133 // CHECK: #pragma omp target teams distribute simd{{$}}
134 // CHECK-NEXT: for (int i = 0; i < 2; ++i)
135 // CHECK-NEXT: a = 2;
136 #pragma omp target teams distribute simd private(argc, b), firstprivate(c, d), collapse(2)
137 for (int i = 0; i < 10; ++i)
138 for (int j = 0; j < 10; ++j)
139 foo();
140 // CHECK: #pragma omp target teams distribute simd private(argc,b) firstprivate(c,d) collapse(2)
141 // CHECK-NEXT: for (int i = 0; i < 10; ++i)
142 // CHECK-NEXT: for (int j = 0; j < 10; ++j)
143 // CHECK-NEXT: foo();
144 for (int i = 0; i < 10; ++i)
145 foo();
146 // CHECK: for (int i = 0; i < 10; ++i)
147 // CHECK-NEXT: foo();
148 #pragma omp target teams distribute simd
149 for (int i = 0; i < 10; ++i)
150 foo();
151 // CHECK: #pragma omp target teams distribute simd
152 // CHECK-NEXT: for (int i = 0; i < 10; ++i)
153 // CHECK-NEXT: foo();
154 #pragma omp target teams distribute simd private(b), firstprivate(argc) shared(d) reduction(+:c) reduction(max:e) num_teams(f) thread_limit(d)
155 for (int k = 0; k < 10; ++k)
156 e += d + argc;
157 // CHECK: #pragma omp target teams distribute simd private(b) firstprivate(argc) shared(d) reduction(+: c) reduction(max: e) num_teams(f) thread_limit(d)
158 // CHECK-NEXT: for (int k = 0; k < 10; ++k)
159 // CHECK-NEXT: e += d + argc;
160 #pragma omp target teams distribute simd simdlen(clen-1)
161 for (int k = 0; k < 10; ++k)
162 e += d + argc;
163 // CHECK: #pragma omp target teams distribute simd simdlen(clen - 1)
164 // CHECK-NEXT: for (int k = 0; k < 10; ++k)
165 // CHECK-NEXT: e += d + argc;
166 #pragma omp target teams distribute simd safelen(clen-1) aligned(arr:alen)
167 for (int k = 0; k < 10; ++k)
168 e += d + argc + arr[k];
169 // CHECK: #pragma omp target teams distribute simd safelen(clen - 1) aligned(arr: alen)
170 // CHECK-NEXT: for (int k = 0; k < 10; ++k)
171 // CHECK-NEXT: e += d + argc + arr[k];
172 return T();
173 }
174
main(int argc,char ** argv)175 int main (int argc, char **argv) {
176 int b = argc, c, d, e, f, g;
177 static int a;
178 // CHECK: static int a;
179 const int clen = 5;
180 const int N = 10;
181 int arr[10];
182 #pragma omp target teams distribute simd
183 for (int i=0; i < 2; ++i)
184 a = 2;
185 // CHECK: #pragma omp target teams distribute simd
186 // CHECK-NEXT: for (int i = 0; i < 2; ++i)
187 // CHECK-NEXT: a = 2;
188 #pragma omp target teams distribute simd private(argc,b),firstprivate(argv, c), collapse(2)
189 for (int i = 0; i < 10; ++i)
190 for (int j = 0; j < 10; ++j)
191 foo();
192 // CHECK: #pragma omp target teams distribute simd private(argc,b) firstprivate(argv,c) collapse(2)
193 // CHECK-NEXT: for (int i = 0; i < 10; ++i)
194 // CHECK-NEXT: for (int j = 0; j < 10; ++j)
195 // CHECK-NEXT: foo();
196 for (int i = 0; i < 10; ++i)
197 foo();
198 // CHECK: for (int i = 0; i < 10; ++i)
199 // CHECK-NEXT: foo();
200 #pragma omp target teams distribute simd
201 for (int i = 0; i < 10; ++i)foo();
202 // CHECK: #pragma omp target teams distribute simd
203 // CHECK-NEXT: for (int i = 0; i < 10; ++i)
204 // CHECK-NEXT: foo();
205 #pragma omp target teams distribute simd private(b), firstprivate(argc) shared(d) reduction(+:c) reduction(max:e) num_teams(f) thread_limit(d)
206 for (int k = 0; k < 10; ++k)
207 e += d + argc;
208 // CHECK: #pragma omp target teams distribute simd private(b) firstprivate(argc) shared(d) reduction(+: c) reduction(max: e) num_teams(f) thread_limit(d)
209 // CHECK-NEXT: for (int k = 0; k < 10; ++k)
210 // CHECK-NEXT: e += d + argc;
211 #pragma omp target teams distribute simd simdlen(clen-1)
212 for (int k = 0; k < 10; ++k)
213 e += d + argc;
214 // CHECK: #pragma omp target teams distribute simd simdlen(clen - 1)
215 // CHECK-NEXT: for (int k = 0; k < 10; ++k)
216 // CHECK-NEXT: e += d + argc;
217 #ifdef OMP51
218 #pragma omp target teams distribute simd safelen(clen-1) aligned(arr:N+6) if(simd:argc) nontemporal(argc, c, d) order(unconstrained:concurrent) allocate(omp_low_lat_mem_alloc:e) firstprivate(e) uses_allocators(omp_low_lat_mem_alloc)
219 #elif OMP5
220 #pragma omp target teams distribute simd safelen(clen-1) aligned(arr:N+6) if(simd:argc) nontemporal(argc, c, d) order(concurrent) allocate(omp_low_lat_mem_alloc:e) firstprivate(e) uses_allocators(omp_low_lat_mem_alloc)
221 #else
222 #pragma omp target teams distribute simd safelen(clen-1) aligned(arr:N+6)
223 #endif // OMP51
224 for (int k = 0; k < 10; ++k)
225 e += d + argc + arr[k];
226 // OMP45: #pragma omp target teams distribute simd safelen(clen - 1) aligned(arr: N + 6)
227 // OMP50: #pragma omp target teams distribute simd safelen(clen - 1) aligned(arr: N + 6) if(simd: argc) nontemporal(argc,c,d) order(concurrent) allocate(omp_low_lat_mem_alloc: e) firstprivate(e) uses_allocators(omp_low_lat_mem_alloc)
228 // OMP51: #pragma omp target teams distribute simd safelen(clen - 1) aligned(arr: N + 6) if(simd: argc) nontemporal(argc,c,d) order(unconstrained: concurrent) allocate(omp_low_lat_mem_alloc: e) firstprivate(e) uses_allocators(omp_low_lat_mem_alloc)
229 // CHECK-NEXT: for (int k = 0; k < 10; ++k)
230 // CHECK-NEXT: e += d + argc + arr[k];
231 return (0);
232 }
233
234 #endif
235