Skip to content

Commit bb97c99

Browse files
erichkeaneerichkeane
committed
[OpenACC] Enable serial/kernels Compute Constructs
So far, all the work we've done for compute constructs has only used 'parallel'. This patch does the work to enable the same logic for 'serial' and 'kernels' constructs as well, since they are the same semantic behavior.
1 parent be3eeea commit bb97c99

File tree

9 files changed

+585
-451
lines changed

9 files changed

+585
-451
lines changed

‎clang/lib/Parse/ParseOpenACC.cpp‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -555,6 +555,8 @@ bool doesDirectiveHaveAssociatedStmt(OpenACCDirectiveKind DirKind) {
555555
default:
556556
return false;
557557
case OpenACCDirectiveKind::Parallel:
558+
case OpenACCDirectiveKind::Serial:
559+
case OpenACCDirectiveKind::Kernels:
558560
return true;
559561
}
560562
llvm_unreachable("Unhandled directive->assoc stmt");
@@ -563,6 +565,8 @@ bool doesDirectiveHaveAssociatedStmt(OpenACCDirectiveKind DirKind) {
563565
unsigned getOpenACCScopeFlags(OpenACCDirectiveKind DirKind) {
564566
switch (DirKind) {
565567
case OpenACCDirectiveKind::Parallel:
568+
case OpenACCDirectiveKind::Serial:
569+
case OpenACCDirectiveKind::Kernels:
566570
// Mark this as a BreakScope/ContinueScope as well as a compute construct
567571
// so that we can diagnose trying to 'break'/'continue' inside of one.
568572
return Scope::BreakScope | Scope::ContinueScope |

‎clang/lib/Sema/SemaOpenACC.cpp‎

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -27,6 +27,8 @@ bool diagnoseConstructAppertainment(Sema &S, OpenACCDirectiveKind K,
2727
// do anything.
2828
break;
2929
case OpenACCDirectiveKind::Parallel:
30+
case OpenACCDirectiveKind::Serial:
31+
case OpenACCDirectiveKind::Kernels:
3032
if (!IsStmt)
3133
return S.Diag(StartLoc, diag::err_acc_construct_appertainment) << K;
3234
break;
@@ -55,6 +57,8 @@ void Sema::ActOnOpenACCConstruct(OpenACCDirectiveKind K,
5557
// rules anywhere.
5658
break;
5759
case OpenACCDirectiveKind::Parallel:
60+
case OpenACCDirectiveKind::Serial:
61+
case OpenACCDirectiveKind::Kernels:
5862
// Nothing to do here, there is no real legalization that needs to happen
5963
// here as these constructs do not take any arguments.
6064
break;
@@ -79,6 +83,8 @@ StmtResult Sema::ActOnEndOpenACCStmtDirective(OpenACCDirectiveKind K,
7983
case OpenACCDirectiveKind::Invalid:
8084
return StmtError();
8185
case OpenACCDirectiveKind::Parallel:
86+
case OpenACCDirectiveKind::Serial:
87+
case OpenACCDirectiveKind::Kernels:
8288
return OpenACCComputeConstruct::Create(
8389
getASTContext(), K, StartLoc, EndLoc,
8490
AssocStmt.isUsable() ? AssocStmt.get() : nullptr);
@@ -92,6 +98,8 @@ StmtResult Sema::ActOnOpenACCAssociatedStmt(OpenACCDirectiveKind K,
9298
default:
9399
llvm_unreachable("Unimplemented associated statement application");
94100
case OpenACCDirectiveKind::Parallel:
101+
case OpenACCDirectiveKind::Serial:
102+
case OpenACCDirectiveKind::Kernels:
95103
// There really isn't any checking here that could happen. As long as we
96104
// have a statement to associate, this should be fine.
97105
// OpenACC 3.3 Section 6:

‎clang/test/ParserOpenACC/parse-clauses.c‎

Lines changed: 404 additions & 437 deletions
Large diffs are not rendered by default.

‎clang/test/ParserOpenACC/parse-constructs.c‎

Lines changed: 8 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -39,23 +39,19 @@ void func() {
3939
// expected-note@+1{{to match this '('}}
4040
#pragma acc parallel( clause list
4141
for(;;){}
42-
// expected-error@+3{{expected clause-list or newline in OpenACC directive}}
43-
// expected-error@+2{{invalid OpenACC clause 'clause'}}
44-
// expected-warning@+1{{OpenACC construct 'serial' not yet implemented, pragma ignored}}
42+
// expected-error@+2{{expected clause-list or newline in OpenACC directive}}
43+
// expected-error@+1{{invalid OpenACC clause 'clause'}}
4544
#pragma acc serial() clause list
4645
for(;;){}
47-
// expected-error@+4{{expected clause-list or newline in OpenACC directive}}
48-
// expected-error@+3{{expected ')'}}
49-
// expected-note@+2{{to match this '('}}
50-
// expected-warning@+1{{OpenACC construct 'serial' not yet implemented, pragma ignored}}
46+
// expected-error@+3{{expected clause-list or newline in OpenACC directive}}
47+
// expected-error@+2{{expected ')'}}
48+
// expected-note@+1{{to match this '('}}
5149
#pragma acc serial( clause list
5250
for(;;){}
53-
// expected-error@+2{{invalid OpenACC clause 'clause'}}
54-
// expected-warning@+1{{OpenACC construct 'serial' not yet implemented, pragma ignored}}
51+
// expected-error@+1{{invalid OpenACC clause 'clause'}}
5552
#pragma acc serial clause list
5653
for(;;){}
57-
// expected-error@+2{{invalid OpenACC clause 'clause'}}
58-
// expected-warning@+1{{OpenACC construct 'kernels' not yet implemented, pragma ignored}}
54+
// expected-error@+1{{invalid OpenACC clause 'clause'}}
5955
#pragma acc kernels clause list
6056
for(;;){}
6157
// expected-error@+2{{invalid OpenACC clause 'clause'}}
@@ -93,8 +89,7 @@ void func() {
9389
// expected-error@+1{{invalid OpenACC clause 'invalid'}}
9490
#pragma acc parallel invalid clause list
9591
for(;;){}
96-
// expected-error@+2{{invalid OpenACC clause 'invalid'}}
97-
// expected-warning@+1{{OpenACC construct 'serial' not yet implemented, pragma ignored}}
92+
// expected-error@+1{{invalid OpenACC clause 'invalid'}}
9893
#pragma acc serial invalid clause list
9994
for(;;){}
10095
// expected-error@+2{{invalid OpenACC clause 'clause'}}

‎clang/test/SemaOpenACC/compute-construct-ast.cpp‎

Lines changed: 50 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -15,6 +15,30 @@ void NormalFunc() {
1515
#pragma acc parallel
1616
{}
1717
}
18+
// FIXME: Add a test once we have clauses for this.
19+
// CHECK-NEXT: OpenACCComputeConstruct {{.*}}serial
20+
// CHECK-NEXT: CompoundStmt
21+
#pragma acc serial
22+
{
23+
#pragma acc serial
24+
// CHECK-NEXT: OpenACCComputeConstruct {{.*}}serial
25+
// CHECK-NEXT: OpenACCComputeConstruct {{.*}}serial
26+
// CHECK-NEXT: CompoundStmt
27+
#pragma acc serial
28+
{}
29+
}
30+
// FIXME: Add a test once we have clauses for this.
31+
// CHECK-NEXT: OpenACCComputeConstruct {{.*}}kernels
32+
// CHECK-NEXT: CompoundStmt
33+
#pragma acc kernels
34+
{
35+
#pragma acc kernels
36+
// CHECK-NEXT: OpenACCComputeConstruct {{.*}}kernels
37+
// CHECK-NEXT: OpenACCComputeConstruct {{.*}}kernels
38+
// CHECK-NEXT: CompoundStmt
39+
#pragma acc kernels
40+
{}
41+
}
1842
}
1943

2044
template<typename T>
@@ -24,6 +48,16 @@ void TemplFunc() {
2448
typename T::type I;
2549
}
2650

51+
#pragma acc serial
52+
{
53+
typename T::type I;
54+
}
55+
56+
#pragma acc kernels
57+
{
58+
typename T::type I;
59+
}
60+
2761
// CHECK-LABEL: FunctionTemplateDecl {{.*}}TemplFunc
2862
// CHECK-NEXT: TemplateTypeParmDecl
2963

@@ -34,6 +68,14 @@ void TemplFunc() {
3468
// CHECK-NEXT: CompoundStmt
3569
// CHECK-NEXT: DeclStmt
3670
// CHECK-NEXT: VarDecl{{.*}} I 'typename T::type'
71+
// CHECK-NEXT: OpenACCComputeConstruct {{.*}}serial
72+
// CHECK-NEXT: CompoundStmt
73+
// CHECK-NEXT: DeclStmt
74+
// CHECK-NEXT: VarDecl{{.*}} I 'typename T::type'
75+
// CHECK-NEXT: OpenACCComputeConstruct {{.*}}kernels
76+
// CHECK-NEXT: CompoundStmt
77+
// CHECK-NEXT: DeclStmt
78+
// CHECK-NEXT: VarDecl{{.*}} I 'typename T::type'
3779

3880
// Check instantiation.
3981
// CHECK-LABEL: FunctionDecl{{.*}} used TemplFunc 'void ()' implicit_instantiation
@@ -45,6 +87,14 @@ void TemplFunc() {
4587
// CHECK-NEXT: CompoundStmt
4688
// CHECK-NEXT: DeclStmt
4789
// CHECK-NEXT: VarDecl{{.*}} I 'typename S::type':'int'
90+
// CHECK-NEXT: OpenACCComputeConstruct {{.*}}serial
91+
// CHECK-NEXT: CompoundStmt
92+
// CHECK-NEXT: DeclStmt
93+
// CHECK-NEXT: VarDecl{{.*}} I 'typename S::type':'int'
94+
// CHECK-NEXT: OpenACCComputeConstruct {{.*}}kernels
95+
// CHECK-NEXT: CompoundStmt
96+
// CHECK-NEXT: DeclStmt
97+
// CHECK-NEXT: VarDecl{{.*}} I 'typename S::type':'int'
4898
}
4999

50100
struct S {

‎clang/test/SemaOpenACC/no-branch-in-out.c‎

Lines changed: 66 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -37,6 +37,18 @@ void BreakContinue() {
3737
break; // expected-error{{invalid branch out of OpenACC Compute Construct}}
3838
}
3939

40+
#pragma acc serial
41+
for(int i = 0; i < 5; ++i) {
42+
if (i > 1)
43+
break; // expected-error{{invalid branch out of OpenACC Compute Construct}}
44+
}
45+
46+
#pragma acc kernels
47+
for(int i = 0; i < 5; ++i) {
48+
if (i > 1)
49+
break; // expected-error{{invalid branch out of OpenACC Compute Construct}}
50+
}
51+
4052
#pragma acc parallel
4153
switch(j) {
4254
case 1:
@@ -99,6 +111,16 @@ void Return() {
99111
return;// expected-error{{invalid return out of OpenACC Compute Construct}}
100112
}
101113

114+
#pragma acc serial
115+
{
116+
return;// expected-error{{invalid return out of OpenACC Compute Construct}}
117+
}
118+
119+
#pragma acc kernels
120+
{
121+
return;// expected-error{{invalid return out of OpenACC Compute Construct}}
122+
}
123+
102124
#pragma acc parallel
103125
{
104126
{
@@ -255,6 +277,34 @@ LABEL13:{}
255277
LABEL14:{}
256278
({goto LABEL14;});
257279
}
280+
281+
282+
283+
({goto LABEL15;});// expected-error{{cannot jump from this goto statement to its label}}
284+
#pragma acc serial// expected-note{{invalid branch into OpenACC Compute Construct}}
285+
{
286+
LABEL15:{}
287+
}
288+
289+
LABEL16:{}
290+
#pragma acc serial// expected-note{{invalid branch out of OpenACC Compute Construct}}
291+
{
292+
({goto LABEL16;});// expected-error{{cannot jump from this goto statement to its label}}
293+
}
294+
295+
296+
({goto LABEL17;});// expected-error{{cannot jump from this goto statement to its label}}
297+
#pragma acc kernels// expected-note{{invalid branch into OpenACC Compute Construct}}
298+
{
299+
LABEL17:{}
300+
}
301+
302+
LABEL18:{}
303+
#pragma acc kernels// expected-note{{invalid branch out of OpenACC Compute Construct}}
304+
{
305+
({goto LABEL18;});// expected-error{{cannot jump from this goto statement to its label}}
306+
}
307+
258308
}
259309

260310
void IndirectGoto1() {
@@ -329,11 +379,27 @@ void DuffsDevice() {
329379
}
330380
}
331381

382+
switch (j) {
383+
#pragma acc kernels
384+
for(int i =0; i < 5; ++i) {
385+
default: // expected-error{{invalid branch into OpenACC Compute Construct}}
386+
{}
387+
}
388+
}
389+
332390
switch (j) {
333391
#pragma acc parallel
334392
for(int i =0; i < 5; ++i) {
335393
case 'a' ... 'z': // expected-error{{invalid branch into OpenACC Compute Construct}}
336394
{}
337395
}
338396
}
397+
398+
switch (j) {
399+
#pragma acc serial
400+
for(int i =0; i < 5; ++i) {
401+
case 'a' ... 'z': // expected-error{{invalid branch into OpenACC Compute Construct}}
402+
{}
403+
}
404+
}
339405
}

‎clang/test/SemaOpenACC/no-branch-in-out.cpp‎

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -147,6 +147,17 @@ void Exceptions() {
147147
throw; // expected-error{{invalid throw out of OpenACC Compute Construct}}
148148
}
149149

150+
#pragma acc serial
151+
for(int i = 0; i < 5; ++i) {
152+
throw; // expected-error{{invalid throw out of OpenACC Compute Construct}}
153+
}
154+
155+
#pragma acc kernels
156+
for(int i = 0; i < 5; ++i) {
157+
throw; // expected-error{{invalid throw out of OpenACC Compute Construct}}
158+
}
159+
160+
150161
#pragma acc parallel
151162
for(int i = 0; i < 5; ++i) {
152163
try {

‎clang/test/SemaOpenACC/parallel-assoc-stmt-inst.cpp‎

Lines changed: 7 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -4,6 +4,10 @@ template<typename T>
44
void Func() {
55
#pragma acc parallel
66
typename T::type I; //#ILOC
7+
#pragma acc serial
8+
typename T::type IS; //#ILOCSERIAL
9+
#pragma acc kernels
10+
typename T::type IK; //#ILOCKERNELS
711
}
812

913
struct S {
@@ -13,6 +17,8 @@ struct S {
1317
void use() {
1418
Func<S>();
1519
// expected-error@#ILOC{{type 'int' cannot be used prior to '::' because it has no members}}
16-
// expected-note@+1{{in instantiation of function template specialization 'Func<int>' requested here}}
20+
// expected-note@+3{{in instantiation of function template specialization 'Func<int>' requested here}}
21+
// expected-error@#ILOCSERIAL{{type 'int' cannot be used prior to '::' because it has no members}}
22+
// expected-error@#ILOCKERNELS{{type 'int' cannot be used prior to '::' because it has no members}}
1723
Func<int>();
1824
}

‎clang/test/SemaOpenACC/parallel-loc-and-stmt.c‎

Lines changed: 27 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3,14 +3,32 @@
33
// expected-error@+1{{OpenACC construct 'parallel' cannot be used here; it can only be used in a statement context}}
44
#pragma acc parallel
55

6+
// expected-error@+1{{OpenACC construct 'serial' cannot be used here; it can only be used in a statement context}}
7+
#pragma acc serial
8+
9+
// expected-error@+1{{OpenACC construct 'kernels' cannot be used here; it can only be used in a statement context}}
10+
#pragma acc kernels
11+
612
// expected-error@+1{{OpenACC construct 'parallel' cannot be used here; it can only be used in a statement context}}
713
#pragma acc parallel
814
int foo;
15+
// expected-error@+1{{OpenACC construct 'serial' cannot be used here; it can only be used in a statement context}}
16+
#pragma acc serial
17+
int foo2;
18+
// expected-error@+1{{OpenACC construct 'kernels' cannot be used here; it can only be used in a statement context}}
19+
#pragma acc kernels
20+
int foo3;
921

1022
struct S {
1123
// expected-error@+1{{OpenACC construct 'parallel' cannot be used here; it can only be used in a statement context}}
1224
#pragma acc parallel
1325
int foo;
26+
// expected-error@+1{{OpenACC construct 'serial' cannot be used here; it can only be used in a statement context}}
27+
#pragma acc serial
28+
int foo2;
29+
// expected-error@+1{{OpenACC construct 'kernels' cannot be used here; it can only be used in a statement context}}
30+
#pragma acc kernels
31+
int foo3;
1432
};
1533

1634
void func() {
@@ -31,6 +49,15 @@ void func() {
3149
#pragma acc parallel
3250
}
3351

52+
{
53+
// expected-error@+2{{expected statement}}
54+
#pragma acc serial
55+
}
56+
{
57+
// expected-error@+2{{expected statement}}
58+
#pragma acc kernels
59+
}
60+
3461
#pragma acc parallel
3562
while(0){}
3663

0 commit comments

Comments
 (0)