Skip to content

Commit c4bfc6f

Browse files
kkwlikkwli
committed
[OpenMP] Support the 'is_device_ptr' clause with 'target parallel for simd' pragma
This patch is to add support of the 'is_device_ptr' clause with the 'target parallel for simd' pragma. Differential Revision: https://reviews.llvm.org/D28402 llvm-svn: 291537
1 parent 4ec493b commit c4bfc6f

File tree

4 files changed

+660
-3
lines changed

4 files changed

+660
-3
lines changed

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

+1-1
Original file line numberDiff line numberDiff line change
@@ -633,7 +633,6 @@ OPENMP_DISTRIBUTE_SIMD_CLAUSE(simdlen)
633633
OPENMP_DISTRIBUTE_SIMD_CLAUSE(reduction)
634634

635635
// Clauses allowed for OpenMP directive 'target parallel for simd'.
636-
// TODO: add target clauses 'is_device_ptr'
637636
OPENMP_TARGET_PARALLEL_FOR_SIMD_CLAUSE(if)
638637
OPENMP_TARGET_PARALLEL_FOR_SIMD_CLAUSE(device)
639638
OPENMP_TARGET_PARALLEL_FOR_SIMD_CLAUSE(map)
@@ -655,6 +654,7 @@ OPENMP_TARGET_PARALLEL_FOR_SIMD_CLAUSE(linear)
655654
OPENMP_TARGET_PARALLEL_FOR_SIMD_CLAUSE(safelen)
656655
OPENMP_TARGET_PARALLEL_FOR_SIMD_CLAUSE(simdlen)
657656
OPENMP_TARGET_PARALLEL_FOR_SIMD_CLAUSE(aligned)
657+
OPENMP_TARGET_PARALLEL_FOR_SIMD_CLAUSE(is_device_ptr)
658658

659659
// Clauses allowed for OpenMP directive 'target simd'.
660660
OPENMP_TARGET_SIMD_CLAUSE(if)

‎clang/lib/Sema/SemaOpenMP.cpp

+4-2
Original file line numberDiff line numberDiff line change
@@ -7451,7 +7451,8 @@ OMPClause *Sema::ActOnOpenMPPrivateClause(ArrayRef<Expr *> VarList,
74517451
CurrDir == OMPD_target_teams ||
74527452
CurrDir == OMPD_target_teams_distribute ||
74537453
CurrDir == OMPD_target_teams_distribute_parallel_for ||
7454-
CurrDir == OMPD_target_teams_distribute_parallel_for_simd) {
7454+
CurrDir == OMPD_target_teams_distribute_parallel_for_simd ||
7455+
CurrDir == OMPD_target_parallel_for_simd) {
74557456
OpenMPClauseKind ConflictKind;
74567457
if (DSAStack->checkMappableExprComponentListsForDecl(
74577458
VD, /*CurrentRegionOnly=*/true,
@@ -7712,7 +7713,8 @@ OMPClause *Sema::ActOnOpenMPFirstprivateClause(ArrayRef<Expr *> VarList,
77127713
CurrDir == OMPD_target_teams ||
77137714
CurrDir == OMPD_target_teams_distribute ||
77147715
CurrDir == OMPD_target_teams_distribute_parallel_for ||
7715-
CurrDir == OMPD_target_teams_distribute_parallel_for_simd) {
7716+
CurrDir == OMPD_target_teams_distribute_parallel_for_simd ||
7717+
CurrDir == OMPD_target_parallel_for_simd) {
77167718
OpenMPClauseKind ConflictKind;
77177719
if (DSAStack->checkMappableExprComponentListsForDecl(
77187720
VD, /*CurrentRegionOnly=*/true,

‎clang/test/OpenMP/target_parallel_for_simd_is_device_ptr_ast_print.cpp

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

0 commit comments

Comments
 (0)