Skip to content

Commit 4101032

Browse files
kkwlikkwli
committed
[OpenMP] Support the 'is_device_ptr' clause with 'target parallel for' pragma
This patch is to add support of the 'is_device_ptr' clause with the 'target parallel for' pragma. Differential Revision: https://reviews.llvm.org/D28255 llvm-svn: 291540
1 parent 17781c7 commit 4101032

File tree

4 files changed

+631
-3
lines changed

4 files changed

+631
-3
lines changed

‎clang/include/clang/Basic/OpenMPKinds.def

+1-1
Original file line numberDiff line numberDiff line change
@@ -483,7 +483,6 @@ OPENMP_TARGET_PARALLEL_CLAUSE(reduction)
483483
OPENMP_TARGET_PARALLEL_CLAUSE(is_device_ptr)
484484

485485
// Clauses allowed for OpenMP directive 'target parallel for'.
486-
// TODO: add target clauses 'is_device_ptr'
487486
OPENMP_TARGET_PARALLEL_FOR_CLAUSE(if)
488487
OPENMP_TARGET_PARALLEL_FOR_CLAUSE(device)
489488
OPENMP_TARGET_PARALLEL_FOR_CLAUSE(map)
@@ -502,6 +501,7 @@ OPENMP_TARGET_PARALLEL_FOR_CLAUSE(collapse)
502501
OPENMP_TARGET_PARALLEL_FOR_CLAUSE(schedule)
503502
OPENMP_TARGET_PARALLEL_FOR_CLAUSE(ordered)
504503
OPENMP_TARGET_PARALLEL_FOR_CLAUSE(linear)
504+
OPENMP_TARGET_PARALLEL_FOR_CLAUSE(is_device_ptr)
505505

506506
// Clauses allowed for OpenMP directive 'target update'.
507507
// TODO More clauses for 'target update' directive.

‎clang/lib/Sema/SemaOpenMP.cpp

+4-2
Original file line numberDiff line numberDiff line change
@@ -7452,7 +7452,8 @@ OMPClause *Sema::ActOnOpenMPPrivateClause(ArrayRef<Expr *> VarList,
74527452
CurrDir == OMPD_target_teams_distribute ||
74537453
CurrDir == OMPD_target_teams_distribute_parallel_for ||
74547454
CurrDir == OMPD_target_teams_distribute_parallel_for_simd ||
7455-
CurrDir == OMPD_target_parallel_for_simd) {
7455+
CurrDir == OMPD_target_parallel_for_simd ||
7456+
CurrDir == OMPD_target_parallel_for) {
74567457
OpenMPClauseKind ConflictKind;
74577458
if (DSAStack->checkMappableExprComponentListsForDecl(
74587459
VD, /*CurrentRegionOnly=*/true,
@@ -7714,7 +7715,8 @@ OMPClause *Sema::ActOnOpenMPFirstprivateClause(ArrayRef<Expr *> VarList,
77147715
CurrDir == OMPD_target_teams_distribute ||
77157716
CurrDir == OMPD_target_teams_distribute_parallel_for ||
77167717
CurrDir == OMPD_target_teams_distribute_parallel_for_simd ||
7717-
CurrDir == OMPD_target_parallel_for_simd) {
7718+
CurrDir == OMPD_target_parallel_for_simd ||
7719+
CurrDir == OMPD_target_parallel_for) {
77187720
OpenMPClauseKind ConflictKind;
77197721
if (DSAStack->checkMappableExprComponentListsForDecl(
77207722
VD, /*CurrentRegionOnly=*/true,

‎clang/test/OpenMP/target_parallel_for_is_device_ptr_ast_print.cpp

+315
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,315 @@
1+
// RUN: %clang_cc1 -verify -fopenmp -std=c++11 -ast-print %s | FileCheck %s
2+
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -emit-pch -o %t %s
3+
// RUN: %clang_cc1 -fopenmp -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
4+
// expected-no-diagnostics
5+
6+
#ifndef HEADER
7+
#define HEADER
8+
9+
void foo() {}
10+
11+
struct ST {
12+
int *a;
13+
};
14+
typedef int arr[10];
15+
typedef ST STarr[10];
16+
struct SA {
17+
const int da[5] = { 0 };
18+
ST g[10];
19+
STarr &rg = g;
20+
int i;
21+
int &j = i;
22+
int *k = &j;
23+
int *&z = k;
24+
int aa[10];
25+
arr &raa = aa;
26+
void func(int arg) {
27+
#pragma omp target parallel for is_device_ptr(k)
28+
for (int i=0; i<100; i++) foo();
29+
30+
#pragma omp target parallel for is_device_ptr(z)
31+
for (int i=0; i<100; i++) foo();
32+
33+
#pragma omp target parallel for is_device_ptr(aa) // OK
34+
for (int i=0; i<100; i++) foo();
35+
36+
#pragma omp target parallel for is_device_ptr(raa) // OK
37+
for (int i=0; i<100; i++) foo();
38+
39+
#pragma omp target parallel for is_device_ptr(g) // OK
40+
for (int i=0; i<100; i++) foo();
41+
42+
#pragma omp target parallel for is_device_ptr(rg) // OK
43+
for (int i=0; i<100; i++) foo();
44+
45+
#pragma omp target parallel for is_device_ptr(da) // OK
46+
for (int i=0; i<100; i++) foo();
47+
48+
return;
49+
}
50+
};
51+
// CHECK: struct SA
52+
// CHECK-NEXT: const int da[5] = {0};
53+
// CHECK-NEXT: ST g[10];
54+
// CHECK-NEXT: STarr &rg = this->g;
55+
// CHECK-NEXT: int i;
56+
// CHECK-NEXT: int &j = this->i;
57+
// CHECK-NEXT: int *k = &this->j;
58+
// CHECK-NEXT: int *&z = this->k;
59+
// CHECK-NEXT: int aa[10];
60+
// CHECK-NEXT: arr &raa = this->aa;
61+
// CHECK-NEXT: func(
62+
// CHECK-NEXT: #pragma omp target parallel for is_device_ptr(this->k)
63+
// CHECK-NEXT: for (int i = 0; i < 100; i++)
64+
// CHECK-NEXT: foo();
65+
// CHECK-NEXT: #pragma omp target parallel for is_device_ptr(this->z)
66+
// CHECK-NEXT: for (int i = 0; i < 100; i++)
67+
// CHECK-NEXT: foo();
68+
// CHECK-NEXT: #pragma omp target parallel for is_device_ptr(this->aa)
69+
// CHECK-NEXT: for (int i = 0; i < 100; i++)
70+
// CHECK-NEXT: foo();
71+
// CHECK-NEXT: #pragma omp target parallel for is_device_ptr(this->raa)
72+
// CHECK-NEXT: for (int i = 0; i < 100; i++)
73+
// CHECK-NEXT: foo();
74+
// CHECK-NEXT: #pragma omp target parallel for is_device_ptr(this->g)
75+
// CHECK-NEXT: for (int i = 0; i < 100; i++)
76+
// CHECK-NEXT: foo();
77+
// CHECK-NEXT: #pragma omp target parallel for is_device_ptr(this->rg)
78+
// CHECK-NEXT: for (int i = 0; i < 100; i++)
79+
// CHECK-NEXT: foo();
80+
// CHECK-NEXT: #pragma omp target parallel for is_device_ptr(this->da)
81+
// CHECK-NEXT: for (int i = 0; i < 100; i++)
82+
// CHECK-NEXT: foo();
83+
84+
struct SB {
85+
unsigned A;
86+
unsigned B;
87+
float Arr[100];
88+
float *Ptr;
89+
float *foo() {
90+
return &Arr[0];
91+
}
92+
};
93+
94+
struct SC {
95+
unsigned A : 2;
96+
unsigned B : 3;
97+
unsigned C;
98+
unsigned D;
99+
float Arr[100];
100+
SB S;
101+
SB ArrS[100];
102+
SB *PtrS;
103+
SB *&RPtrS;
104+
float *Ptr;
105+
106+
SC(SB *&_RPtrS) : RPtrS(_RPtrS) {}
107+
};
108+
109+
union SD {
110+
unsigned A;
111+
float B;
112+
};
113+
114+
struct S1;
115+
extern S1 a;
116+
class S2 {
117+
mutable int a;
118+
public:
119+
S2():a(0) { }
120+
S2(S2 &s2):a(s2.a) { }
121+
static float S2s;
122+
static const float S2sc;
123+
};
124+
const float S2::S2sc = 0;
125+
const S2 b;
126+
const S2 ba[5];
127+
class S3 {
128+
int a;
129+
public:
130+
S3():a(0) { }
131+
S3(S3 &s3):a(s3.a) { }
132+
};
133+
const S3 c;
134+
const S3 ca[5];
135+
extern const int f;
136+
class S4 {
137+
int a;
138+
S4();
139+
S4(const S4 &s4);
140+
public:
141+
S4(int v):a(v) { }
142+
};
143+
class S5 {
144+
int a;
145+
S5():a(0) {}
146+
S5(const S5 &s5):a(s5.a) { }
147+
public:
148+
S5(int v):a(v) { }
149+
};
150+
151+
S3 h;
152+
#pragma omp threadprivate(h)
153+
154+
typedef struct {
155+
int a;
156+
} S6;
157+
158+
template <typename T>
159+
T tmain(T argc) {
160+
const T da[5] = { 0 };
161+
S6 h[10];
162+
auto &rh = h;
163+
T i;
164+
T &j = i;
165+
T *k = &j;
166+
T *&z = k;
167+
T aa[10];
168+
auto &raa = aa;
169+
#pragma omp target parallel for is_device_ptr(k)
170+
for (int i=0; i<100; i++) foo();
171+
#pragma omp target parallel for is_device_ptr(z)
172+
for (int i=0; i<100; i++) foo();
173+
#pragma omp target parallel for is_device_ptr(aa)
174+
for (int i=0; i<100; i++) foo();
175+
#pragma omp target parallel for is_device_ptr(raa)
176+
for (int i=0; i<100; i++) foo();
177+
#pragma omp target parallel for is_device_ptr(h)
178+
for (int i=0; i<100; i++) foo();
179+
#pragma omp target parallel for is_device_ptr(rh)
180+
for (int i=0; i<100; i++) foo();
181+
#pragma omp target parallel for is_device_ptr(da)
182+
for (int i=0; i<100; i++) foo();
183+
return 0;
184+
}
185+
186+
// CHECK: template<> int tmain<int>(int argc) {
187+
// CHECK-NEXT: const int da[5] = {0};
188+
// CHECK-NEXT: S6 h[10];
189+
// CHECK-NEXT: auto &rh = h;
190+
// CHECK-NEXT: int i;
191+
// CHECK-NEXT: int &j = i;
192+
// CHECK-NEXT: int *k = &j;
193+
// CHECK-NEXT: int *&z = k;
194+
// CHECK-NEXT: int aa[10];
195+
// CHECK-NEXT: auto &raa = aa;
196+
// CHECK-NEXT: #pragma omp target parallel for is_device_ptr(k)
197+
// CHECK-NEXT: for (int i = 0; i < 100; i++)
198+
// CHECK-NEXT: foo();
199+
// CHECK-NEXT: #pragma omp target parallel for is_device_ptr(z)
200+
// CHECK-NEXT: for (int i = 0; i < 100; i++)
201+
// CHECK-NEXT: foo();
202+
// CHECK-NEXT: #pragma omp target parallel for is_device_ptr(aa)
203+
// CHECK-NEXT: for (int i = 0; i < 100; i++)
204+
// CHECK-NEXT: foo();
205+
// CHECK-NEXT: #pragma omp target parallel for is_device_ptr(raa)
206+
// CHECK-NEXT: for (int i = 0; i < 100; i++)
207+
// CHECK-NEXT: foo();
208+
// CHECK-NEXT: #pragma omp target parallel for is_device_ptr(h)
209+
// CHECK-NEXT: for (int i = 0; i < 100; i++)
210+
// CHECK-NEXT: foo();
211+
// CHECK-NEXT: #pragma omp target parallel for is_device_ptr(rh)
212+
// CHECK-NEXT: for (int i = 0; i < 100; i++)
213+
// CHECK-NEXT: foo();
214+
// CHECK-NEXT: #pragma omp target parallel for is_device_ptr(da)
215+
// CHECK-NEXT: for (int i = 0; i < 100; i++)
216+
// CHECK-NEXT: foo();
217+
218+
// CHECK: template<> int *tmain<int *>(int *argc) {
219+
// CHECK-NEXT: int *const da[5] = {0};
220+
// CHECK-NEXT: S6 h[10];
221+
// CHECK-NEXT: auto &rh = h;
222+
// CHECK-NEXT: int *i;
223+
// CHECK-NEXT: int *&j = i;
224+
// CHECK-NEXT: int **k = &j;
225+
// CHECK-NEXT: int **&z = k;
226+
// CHECK-NEXT: int *aa[10];
227+
// CHECK-NEXT: auto &raa = aa;
228+
// CHECK-NEXT: #pragma omp target parallel for is_device_ptr(k)
229+
// CHECK-NEXT: for (int i = 0; i < 100; i++)
230+
// CHECK-NEXT: foo();
231+
// CHECK-NEXT: #pragma omp target parallel for is_device_ptr(z)
232+
// CHECK-NEXT: for (int i = 0; i < 100; i++)
233+
// CHECK-NEXT: foo();
234+
// CHECK-NEXT: #pragma omp target parallel for is_device_ptr(aa)
235+
// CHECK-NEXT: for (int i = 0; i < 100; i++)
236+
// CHECK-NEXT: foo();
237+
// CHECK-NEXT: #pragma omp target parallel for is_device_ptr(raa)
238+
// CHECK-NEXT: for (int i = 0; i < 100; i++)
239+
// CHECK-NEXT: foo();
240+
// CHECK-NEXT: #pragma omp target parallel for is_device_ptr(h)
241+
// CHECK-NEXT: for (int i = 0; i < 100; i++)
242+
// CHECK-NEXT: foo();
243+
// CHECK-NEXT: #pragma omp target parallel for is_device_ptr(rh)
244+
// CHECK-NEXT: for (int i = 0; i < 100; i++)
245+
// CHECK-NEXT: foo();
246+
// CHECK-NEXT: #pragma omp target parallel for is_device_ptr(da)
247+
// CHECK-NEXT: for (int i = 0; i < 100; i++)
248+
// CHECK-NEXT: foo();
249+
250+
// CHECK-LABEL: int main(int argc, char **argv) {
251+
int main(int argc, char **argv) {
252+
const int da[5] = { 0 };
253+
S6 h[10];
254+
auto &rh = h;
255+
int i;
256+
int &j = i;
257+
int *k = &j;
258+
int *&z = k;
259+
int aa[10];
260+
auto &raa = aa;
261+
// CHECK-NEXT: const int da[5] = {0};
262+
// CHECK-NEXT: S6 h[10];
263+
// CHECK-NEXT: auto &rh = h;
264+
// CHECK-NEXT: int i;
265+
// CHECK-NEXT: int &j = i;
266+
// CHECK-NEXT: int *k = &j;
267+
// CHECK-NEXT: int *&z = k;
268+
// CHECK-NEXT: int aa[10];
269+
// CHECK-NEXT: auto &raa = aa;
270+
#pragma omp target parallel for is_device_ptr(k)
271+
for (int i=0; i<100; i++) foo();
272+
// CHECK-NEXT: #pragma omp target parallel for is_device_ptr(k)
273+
// CHECK-NEXT: for (int i = 0; i < 100; i++)
274+
// CHECK-NEXT: foo();
275+
276+
#pragma omp target parallel for is_device_ptr(z)
277+
for (int i=0; i<100; i++) foo();
278+
// CHECK-NEXT: #pragma omp target parallel for is_device_ptr(z)
279+
// CHECK-NEXT: for (int i = 0; i < 100; i++)
280+
// CHECK-NEXT: foo();
281+
282+
#pragma omp target parallel for is_device_ptr(aa)
283+
for (int i=0; i<100; i++) foo();
284+
// CHECK-NEXT: #pragma omp target parallel for is_device_ptr(aa)
285+
// CHECK-NEXT: for (int i = 0; i < 100; i++)
286+
// CHECK-NEXT: foo();
287+
288+
#pragma omp target parallel for is_device_ptr(raa)
289+
for (int i=0; i<100; i++) foo();
290+
// CHECK-NEXT: #pragma omp target parallel for is_device_ptr(raa)
291+
// CHECK-NEXT: for (int i = 0; i < 100; i++)
292+
// CHECK-NEXT: foo();
293+
294+
#pragma omp target parallel for is_device_ptr(h)
295+
for (int i=0; i<100; i++) foo();
296+
// CHECK-NEXT: #pragma omp target parallel for is_device_ptr(h)
297+
// CHECK-NEXT: for (int i = 0; i < 100; i++)
298+
// CHECK-NEXT: foo();
299+
300+
#pragma omp target parallel for is_device_ptr(rh)
301+
for (int i=0; i<100; i++) foo();
302+
// CHECK-NEXT: #pragma omp target parallel for is_device_ptr(rh)
303+
// CHECK-NEXT: for (int i = 0; i < 100; i++)
304+
// CHECK-NEXT: foo();
305+
306+
#pragma omp target parallel for is_device_ptr(da)
307+
for (int i=0; i<100; i++) foo();
308+
// CHECK-NEXT: #pragma omp target parallel for is_device_ptr(da)
309+
// CHECK-NEXT: for (int i = 0; i < 100; i++)
310+
// CHECK-NEXT: foo();
311+
return tmain<int>(argc) + *tmain<int *>(&argc);
312+
}
313+
314+
315+
#endif

0 commit comments

Comments
 (0)