Skip to content

Commit b52fe2d

Browse files
erichkeaneerichkeane
committed
[OpenACC] Implement 'gang' clause parsing.
The 'gang' clause takes a 'gang-arg-list', which is one of three 'tag' values, followed by either an 'int-expr' or a 'size-expr', both of which we already have parsing functions for. The optional tag values are only slightly complicated, as one is a keyword (static), so mild modifications needed to be made for that.
1 parent 21d75ee commit b52fe2d

File tree

4 files changed

+228
-1
lines changed

4 files changed

+228
-1
lines changed

‎clang/include/clang/Basic/OpenACCKinds.h

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -243,6 +243,8 @@ enum class OpenACCClauseKind {
243243
Async,
244244
/// 'tile' clause, allowed on 'loop' and Combined constructs.
245245
Tile,
246+
/// 'gang' clause, allowed on 'loop' and Combined constructs.
247+
Gang,
246248

247249
/// Represents an invalid clause, for the purposes of parsing.
248250
Invalid,
@@ -371,6 +373,9 @@ inline const StreamingDiagnostic &operator<<(const StreamingDiagnostic &Out,
371373
case OpenACCClauseKind::Tile:
372374
return Out << "tile";
373375

376+
case OpenACCClauseKind::Gang:
377+
return Out << "gang";
378+
374379
case OpenACCClauseKind::Invalid:
375380
return Out << "<invalid>";
376381
}

‎clang/include/clang/Parse/Parser.h

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3605,6 +3605,10 @@ class Parser : public CodeCompletionHandler {
36053605
bool ParseOpenACCSizeExpr();
36063606
/// Parses a comma delimited list of 'size-expr's.
36073607
bool ParseOpenACCSizeExprList();
3608+
/// Parses a 'gang-arg-list', used for the 'gang' clause.
3609+
bool ParseOpenACCGangArgList();
3610+
/// Parses a 'gang-arg', used for the 'gang' clause.
3611+
bool ParseOpenACCGangArg();
36083612

36093613
private:
36103614
//===--------------------------------------------------------------------===//

‎clang/lib/Parse/ParseOpenACC.cpp

Lines changed: 69 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -110,6 +110,7 @@ OpenACCClauseKind getOpenACCClauseKind(Token Tok) {
110110
.Case("dtype", OpenACCClauseKind::DType)
111111
.Case("finalize", OpenACCClauseKind::Finalize)
112112
.Case("firstprivate", OpenACCClauseKind::FirstPrivate)
113+
.Case("gang", OpenACCClauseKind::Gang)
113114
.Case("host", OpenACCClauseKind::Host)
114115
.Case("if", OpenACCClauseKind::If)
115116
.Case("if_present", OpenACCClauseKind::IfPresent)
@@ -165,9 +166,14 @@ enum class OpenACCSpecialTokenKind {
165166
Force,
166167
Num,
167168
Length,
169+
Dim,
170+
Static,
168171
};
169172

170173
bool isOpenACCSpecialToken(OpenACCSpecialTokenKind Kind, Token Tok) {
174+
if (Tok.is(tok::kw_static) && Kind == OpenACCSpecialTokenKind::Static)
175+
return true;
176+
171177
if (!Tok.is(tok::identifier))
172178
return false;
173179

@@ -186,6 +192,10 @@ bool isOpenACCSpecialToken(OpenACCSpecialTokenKind Kind, Token Tok) {
186192
return Tok.getIdentifierInfo()->isStr("num");
187193
case OpenACCSpecialTokenKind::Length:
188194
return Tok.getIdentifierInfo()->isStr("length");
195+
case OpenACCSpecialTokenKind::Dim:
196+
return Tok.getIdentifierInfo()->isStr("dim");
197+
case OpenACCSpecialTokenKind::Static:
198+
return Tok.getIdentifierInfo()->isStr("static");
189199
}
190200
llvm_unreachable("Unknown 'Kind' Passed");
191201
}
@@ -464,6 +474,7 @@ ClauseParensKind getClauseParensKind(OpenACCDirectiveKind DirKind,
464474
case OpenACCClauseKind::Async:
465475
case OpenACCClauseKind::Worker:
466476
case OpenACCClauseKind::Vector:
477+
case OpenACCClauseKind::Gang:
467478
return ClauseParensKind::Optional;
468479

469480
case OpenACCClauseKind::Default:
@@ -632,7 +643,8 @@ bool Parser::ParseOpenACCSizeExpr() {
632643
// The size-expr ends up being ambiguous when only looking at the current
633644
// token, as it could be a deref of a variable/expression.
634645
if (getCurToken().is(tok::star) &&
635-
NextToken().isOneOf(tok::comma, tok::r_paren)) {
646+
NextToken().isOneOf(tok::comma, tok::r_paren,
647+
tok::annot_pragma_openacc_end)) {
636648
ConsumeToken();
637649
return false;
638650
}
@@ -661,6 +673,58 @@ bool Parser::ParseOpenACCSizeExprList() {
661673
return false;
662674
}
663675

676+
/// OpenACC 3.3 Section 2.9:
677+
///
678+
/// where gang-arg is one of:
679+
/// [num:]int-expr
680+
/// dim:int-expr
681+
/// static:size-expr
682+
bool Parser::ParseOpenACCGangArg() {
683+
684+
if (isOpenACCSpecialToken(OpenACCSpecialTokenKind::Static, getCurToken()) &&
685+
NextToken().is(tok::colon)) {
686+
// 'static' just takes a size-expr, which is an int-expr or an asterisk.
687+
ConsumeToken();
688+
ConsumeToken();
689+
return ParseOpenACCSizeExpr();
690+
}
691+
692+
if (isOpenACCSpecialToken(OpenACCSpecialTokenKind::Dim, getCurToken()) &&
693+
NextToken().is(tok::colon)) {
694+
ConsumeToken();
695+
ConsumeToken();
696+
return ParseOpenACCIntExpr().isInvalid();
697+
}
698+
699+
if (isOpenACCSpecialToken(OpenACCSpecialTokenKind::Num, getCurToken()) &&
700+
NextToken().is(tok::colon)) {
701+
ConsumeToken();
702+
ConsumeToken();
703+
// Fallthrough to the 'int-expr' handling for when 'num' is omitted.
704+
}
705+
// This is just the 'num' case where 'num' is optional.
706+
return ParseOpenACCIntExpr().isInvalid();
707+
}
708+
709+
bool Parser::ParseOpenACCGangArgList() {
710+
if (ParseOpenACCGangArg()) {
711+
SkipUntil(tok::r_paren, tok::annot_pragma_openacc_end,
712+
Parser::StopBeforeMatch);
713+
return false;
714+
}
715+
716+
while (!getCurToken().isOneOf(tok::r_paren, tok::annot_pragma_openacc_end)) {
717+
ExpectAndConsume(tok::comma);
718+
719+
if (ParseOpenACCGangArg()) {
720+
SkipUntil(tok::r_paren, tok::annot_pragma_openacc_end,
721+
Parser::StopBeforeMatch);
722+
return false;
723+
}
724+
}
725+
return false;
726+
}
727+
664728
// The OpenACC Clause List is a comma or space-delimited list of clauses (see
665729
// the comment on ParseOpenACCClauseList). The concept of a 'clause' doesn't
666730
// really have its owner grammar and each individual one has its own definition.
@@ -839,6 +903,10 @@ bool Parser::ParseOpenACCClauseParams(OpenACCDirectiveKind DirKind,
839903
return true;
840904
break;
841905
}
906+
case OpenACCClauseKind::Gang:
907+
if (ParseOpenACCGangArgList())
908+
return true;
909+
break;
842910
default:
843911
llvm_unreachable("Not an optional parens type?");
844912
}

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

Lines changed: 150 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1111,6 +1111,156 @@ void Tile() {
11111111
for(;;){}
11121112
}
11131113

1114+
void Gang() {
1115+
// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
1116+
#pragma acc loop gang
1117+
for(;;){}
1118+
// expected-error@+4{{expected expression}}
1119+
// expected-error@+3{{expected ')'}}
1120+
// expected-note@+2{{to match this '('}}
1121+
// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
1122+
#pragma acc loop gang(
1123+
for(;;){}
1124+
// expected-error@+2{{expected expression}}
1125+
// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
1126+
#pragma acc loop gang()
1127+
for(;;){}
1128+
1129+
// expected-error@+2{{expected expression}}
1130+
// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
1131+
#pragma acc loop gang(5, *)
1132+
for(;;){}
1133+
1134+
// expected-error@+2{{expected expression}}
1135+
// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
1136+
#pragma acc loop gang(*)
1137+
for(;;){}
1138+
1139+
// expected-error@+2{{expected expression}}
1140+
// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
1141+
#pragma acc loop gang(5, num:*)
1142+
for(;;){}
1143+
1144+
// expected-error@+2{{expected expression}}
1145+
// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
1146+
#pragma acc loop gang(num:5, *)
1147+
for(;;){}
1148+
1149+
// expected-error@+2{{expected expression}}
1150+
// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
1151+
#pragma acc loop gang(num:5, num:*)
1152+
for(;;){}
1153+
1154+
// expected-error@+2{{expected expression}}
1155+
// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
1156+
#pragma acc loop gang(num:*)
1157+
for(;;){}
1158+
1159+
// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
1160+
#pragma acc loop gang(dim:5)
1161+
for(;;){}
1162+
1163+
// expected-error@+2{{expected expression}}
1164+
// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
1165+
#pragma acc loop gang(dim:5, dim:*)
1166+
for(;;){}
1167+
1168+
// expected-error@+2{{expected expression}}
1169+
// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
1170+
#pragma acc loop gang(dim:*)
1171+
for(;;){}
1172+
1173+
// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
1174+
#pragma acc loop gang(static:*)
1175+
for(;;){}
1176+
1177+
// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
1178+
#pragma acc loop gang(static:*, static:5)
1179+
for(;;){}
1180+
1181+
// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
1182+
#pragma acc loop gang(static:*, 5)
1183+
for(;;){}
1184+
1185+
// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
1186+
#pragma acc loop gang(static:45, 5)
1187+
for(;;){}
1188+
1189+
// expected-error@+4{{expected expression}}
1190+
// expected-error@+3{{expected ')'}}
1191+
// expected-note@+2{{to match this '('}}
1192+
// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
1193+
#pragma acc loop gang(static:45,
1194+
for(;;){}
1195+
1196+
// expected-error@+3{{expected ')'}}
1197+
// expected-note@+2{{to match this '('}}
1198+
// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
1199+
#pragma acc loop gang(static:45
1200+
for(;;){}
1201+
1202+
// expected-error@+4{{expected expression}}
1203+
// expected-error@+3{{expected ')'}}
1204+
// expected-note@+2{{to match this '('}}
1205+
// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
1206+
#pragma acc loop gang(static:*,
1207+
for(;;){}
1208+
1209+
// expected-error@+3{{expected ')'}}
1210+
// expected-note@+2{{to match this '('}}
1211+
// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
1212+
#pragma acc loop gang(static:*
1213+
for(;;){}
1214+
1215+
// expected-error@+4{{expected expression}}
1216+
// expected-error@+3{{expected ')'}}
1217+
// expected-note@+2{{to match this '('}}
1218+
// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
1219+
#pragma acc loop gang(45,
1220+
for(;;){}
1221+
1222+
// expected-error@+3{{expected ')'}}
1223+
// expected-note@+2{{to match this '('}}
1224+
// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
1225+
#pragma acc loop gang(45
1226+
for(;;){}
1227+
1228+
// expected-error@+4{{expected expression}}
1229+
// expected-error@+3{{expected ')'}}
1230+
// expected-note@+2{{to match this '('}}
1231+
// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
1232+
#pragma acc loop gang(num:45,
1233+
for(;;){}
1234+
1235+
// expected-error@+3{{expected ')'}}
1236+
// expected-note@+2{{to match this '('}}
1237+
// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
1238+
#pragma acc loop gang(num:45
1239+
for(;;){}
1240+
1241+
// expected-error@+4{{expected expression}}
1242+
// expected-error@+3{{expected ')'}}
1243+
// expected-note@+2{{to match this '('}}
1244+
// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
1245+
#pragma acc loop gang(dim:45,
1246+
for(;;){}
1247+
1248+
// expected-error@+3{{expected ')'}}
1249+
// expected-note@+2{{to match this '('}}
1250+
// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
1251+
#pragma acc loop gang(dim:45
1252+
for(;;){}
1253+
1254+
// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
1255+
#pragma acc loop gang(static:*, dim:returns_int(), 5)
1256+
for(;;){}
1257+
1258+
// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
1259+
#pragma acc loop gang(num: 32, static:*, dim:returns_int(), 5)
1260+
for(;;){}
1261+
1262+
}
1263+
11141264
// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
11151265
#pragma acc routine worker, vector, seq, nohost
11161266
void bar();

0 commit comments

Comments
 (0)