Skip to content

Commit 30cfe2b

Browse files
erichkeaneerichkeane
committed
[OpenACC] Implement 'async' clause sema for compute constructs
This is a pretty simple clause, it takes an 'async-argument', which effectively needs to be just parsed as an 'int' argument, since it can be an arbitrarly integer at runtime (and negative values are legal for implementation defined values). This patch also cleans up the async-argument parsing, so 'wait' got some minor quality-of-life improvements for parsing (both clause and construct).
1 parent f548c4d commit 30cfe2b

21 files changed

+461
-27
lines changed

‎clang/include/clang/AST/OpenACCClause.h

Lines changed: 13 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -227,7 +227,8 @@ class OpenACCClauseWithSingleIntExpr : public OpenACCClauseWithExprs {
227227
SourceLocation EndLoc)
228228
: OpenACCClauseWithExprs(K, BeginLoc, LParenLoc, EndLoc),
229229
IntExpr(IntExpr) {
230-
setExprs(MutableArrayRef<Expr *>{&this->IntExpr, 1});
230+
if (IntExpr)
231+
setExprs(MutableArrayRef<Expr *>{&this->IntExpr, 1});
231232
}
232233

233234
public:
@@ -260,6 +261,17 @@ class OpenACCVectorLengthClause : public OpenACCClauseWithSingleIntExpr {
260261
Expr *IntExpr, SourceLocation EndLoc);
261262
};
262263

264+
class OpenACCAsyncClause : public OpenACCClauseWithSingleIntExpr {
265+
OpenACCAsyncClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
266+
Expr *IntExpr, SourceLocation EndLoc);
267+
268+
public:
269+
static OpenACCAsyncClause *Create(const ASTContext &C,
270+
SourceLocation BeginLoc,
271+
SourceLocation LParenLoc, Expr *IntExpr,
272+
SourceLocation EndLoc);
273+
};
274+
263275
/// Represents a clause with one or more 'var' objects, represented as an expr,
264276
/// as its arguments. Var-list is expected to be stored in trailing storage.
265277
/// For now, we're just storing the original expression in its entirety, unlike

‎clang/include/clang/Basic/OpenACCClauses.def

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -21,6 +21,7 @@
2121
#define CLAUSE_ALIAS(ALIAS_NAME, CLAUSE_NAME)
2222
#endif
2323

24+
VISIT_CLAUSE(Async)
2425
VISIT_CLAUSE(Attach)
2526
VISIT_CLAUSE(Copy)
2627
CLAUSE_ALIAS(PCopy, Copy)

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

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3698,7 +3698,9 @@ class Parser : public CodeCompletionHandler {
36983698
bool ParseOpenACCDeviceTypeList();
36993699
/// Parses the 'async-argument', which is an integral value with two
37003700
/// 'special' values that are likely negative (but come from Macros).
3701-
ExprResult ParseOpenACCAsyncArgument();
3701+
OpenACCIntExprParseResult ParseOpenACCAsyncArgument(OpenACCDirectiveKind DK,
3702+
OpenACCClauseKind CK,
3703+
SourceLocation Loc);
37023704
/// Parses the 'size-expr', which is an integral value, or an asterisk.
37033705
bool ParseOpenACCSizeExpr();
37043706
/// Parses a comma delimited list of 'size-expr's.

‎clang/include/clang/Sema/SemaOpenACC.h

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -101,16 +101,24 @@ class SemaOpenACC : public SemaBase {
101101
unsigned getNumIntExprs() const {
102102
assert((ClauseKind == OpenACCClauseKind::NumGangs ||
103103
ClauseKind == OpenACCClauseKind::NumWorkers ||
104+
ClauseKind == OpenACCClauseKind::Async ||
104105
ClauseKind == OpenACCClauseKind::VectorLength) &&
105106
"Parsed clause kind does not have a int exprs");
107+
//
108+
// 'async' has an optional IntExpr, so be tolerant of that.
109+
if (ClauseKind == OpenACCClauseKind::Async &&
110+
std::holds_alternative<std::monostate>(Details))
111+
return 0;
106112
return std::get<IntExprDetails>(Details).IntExprs.size();
107113
}
108114

109115
ArrayRef<Expr *> getIntExprs() {
110116
assert((ClauseKind == OpenACCClauseKind::NumGangs ||
111117
ClauseKind == OpenACCClauseKind::NumWorkers ||
118+
ClauseKind == OpenACCClauseKind::Async ||
112119
ClauseKind == OpenACCClauseKind::VectorLength) &&
113120
"Parsed clause kind does not have a int exprs");
121+
114122
return std::get<IntExprDetails>(Details).IntExprs;
115123
}
116124

@@ -190,13 +198,15 @@ class SemaOpenACC : public SemaBase {
190198
void setIntExprDetails(ArrayRef<Expr *> IntExprs) {
191199
assert((ClauseKind == OpenACCClauseKind::NumGangs ||
192200
ClauseKind == OpenACCClauseKind::NumWorkers ||
201+
ClauseKind == OpenACCClauseKind::Async ||
193202
ClauseKind == OpenACCClauseKind::VectorLength) &&
194203
"Parsed clause kind does not have a int exprs");
195204
Details = IntExprDetails{{IntExprs.begin(), IntExprs.end()}};
196205
}
197206
void setIntExprDetails(llvm::SmallVector<Expr *> &&IntExprs) {
198207
assert((ClauseKind == OpenACCClauseKind::NumGangs ||
199208
ClauseKind == OpenACCClauseKind::NumWorkers ||
209+
ClauseKind == OpenACCClauseKind::Async ||
200210
ClauseKind == OpenACCClauseKind::VectorLength) &&
201211
"Parsed clause kind does not have a int exprs");
202212
Details = IntExprDetails{std::move(IntExprs)};

‎clang/lib/AST/OpenACCClause.cpp

Lines changed: 29 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -127,6 +127,26 @@ OpenACCVectorLengthClause::Create(const ASTContext &C, SourceLocation BeginLoc,
127127
OpenACCVectorLengthClause(BeginLoc, LParenLoc, IntExpr, EndLoc);
128128
}
129129

130+
OpenACCAsyncClause::OpenACCAsyncClause(SourceLocation BeginLoc,
131+
SourceLocation LParenLoc, Expr *IntExpr,
132+
SourceLocation EndLoc)
133+
: OpenACCClauseWithSingleIntExpr(OpenACCClauseKind::Async, BeginLoc,
134+
LParenLoc, IntExpr, EndLoc) {
135+
assert((!IntExpr || IntExpr->isInstantiationDependent() ||
136+
IntExpr->getType()->isIntegerType()) &&
137+
"Condition expression type not scalar/dependent");
138+
}
139+
140+
OpenACCAsyncClause *OpenACCAsyncClause::Create(const ASTContext &C,
141+
SourceLocation BeginLoc,
142+
SourceLocation LParenLoc,
143+
Expr *IntExpr,
144+
SourceLocation EndLoc) {
145+
void *Mem =
146+
C.Allocate(sizeof(OpenACCAsyncClause), alignof(OpenACCAsyncClause));
147+
return new (Mem) OpenACCAsyncClause(BeginLoc, LParenLoc, IntExpr, EndLoc);
148+
}
149+
130150
OpenACCNumGangsClause *OpenACCNumGangsClause::Create(const ASTContext &C,
131151
SourceLocation BeginLoc,
132152
SourceLocation LParenLoc,
@@ -287,6 +307,15 @@ void OpenACCClausePrinter::VisitVectorLengthClause(
287307
OS << ")";
288308
}
289309

310+
void OpenACCClausePrinter::VisitAsyncClause(const OpenACCAsyncClause &C) {
311+
OS << "async";
312+
if (C.hasIntExpr()) {
313+
OS << "(";
314+
printExpr(C.getIntExpr());
315+
OS << ")";
316+
}
317+
}
318+
290319
void OpenACCClausePrinter::VisitPrivateClause(const OpenACCPrivateClause &C) {
291320
OS << "private(";
292321
llvm::interleaveComma(C.getVarList(), OS,

‎clang/lib/AST/StmtProfile.cpp

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2573,6 +2573,11 @@ void OpenACCClauseProfiler::VisitVectorLengthClause(
25732573
"vector_length clause requires a valid int expr");
25742574
Profiler.VisitStmt(Clause.getIntExpr());
25752575
}
2576+
2577+
void OpenACCClauseProfiler::VisitAsyncClause(const OpenACCAsyncClause &Clause) {
2578+
if (Clause.hasIntExpr())
2579+
Profiler.VisitStmt(Clause.getIntExpr());
2580+
}
25762581
} // namespace
25772582

25782583
void StmtProfiler::VisitOpenACCComputeConstruct(

‎clang/lib/AST/TextNodeDumper.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -397,6 +397,7 @@ void TextNodeDumper::Visit(const OpenACCClause *C) {
397397
case OpenACCClauseKind::Default:
398398
OS << '(' << cast<OpenACCDefaultClause>(C)->getDefaultClauseKind() << ')';
399399
break;
400+
case OpenACCClauseKind::Async:
400401
case OpenACCClauseKind::Attach:
401402
case OpenACCClauseKind::Copy:
402403
case OpenACCClauseKind::PCopy:

‎clang/lib/Parse/ParseOpenACC.cpp

Lines changed: 22 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -1081,7 +1081,12 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
10811081
break;
10821082
}
10831083
case OpenACCClauseKind::Async: {
1084-
ExprResult AsyncArg = ParseOpenACCAsyncArgument();
1084+
ExprResult AsyncArg =
1085+
ParseOpenACCAsyncArgument(OpenACCDirectiveKind::Invalid,
1086+
OpenACCClauseKind::Async, ClauseLoc)
1087+
.first;
1088+
ParsedClause.setIntExprDetails(AsyncArg.isUsable() ? AsyncArg.get()
1089+
: nullptr);
10851090
if (AsyncArg.isInvalid()) {
10861091
Parens.skipToEnd();
10871092
return OpenACCCanContinue();
@@ -1120,8 +1125,10 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
11201125
/// defined in the C header file and the Fortran openacc module. The special
11211126
/// values are negative values, so as not to conflict with a user-specified
11221127
/// nonnegative async-argument.
1123-
ExprResult Parser::ParseOpenACCAsyncArgument() {
1124-
return getActions().CorrectDelayedTyposInExpr(ParseAssignmentExpression());
1128+
Parser::OpenACCIntExprParseResult
1129+
Parser::ParseOpenACCAsyncArgument(OpenACCDirectiveKind DK, OpenACCClauseKind CK,
1130+
SourceLocation Loc) {
1131+
return ParseOpenACCIntExpr(DK, CK, Loc);
11251132
}
11261133

11271134
/// OpenACC 3.3, section 2.16:
@@ -1137,14 +1144,12 @@ bool Parser::ParseOpenACCWaitArgument(SourceLocation Loc, bool IsDirective) {
11371144
// Consume colon.
11381145
ConsumeToken();
11391146

1140-
ExprResult IntExpr =
1141-
ParseOpenACCIntExpr(IsDirective ? OpenACCDirectiveKind::Wait
1142-
: OpenACCDirectiveKind::Invalid,
1143-
IsDirective ? OpenACCClauseKind::Invalid
1144-
: OpenACCClauseKind::Wait,
1145-
Loc)
1146-
.first;
1147-
if (IntExpr.isInvalid())
1147+
OpenACCIntExprParseResult Res = ParseOpenACCIntExpr(
1148+
IsDirective ? OpenACCDirectiveKind::Wait
1149+
: OpenACCDirectiveKind::Invalid,
1150+
IsDirective ? OpenACCClauseKind::Invalid : OpenACCClauseKind::Wait,
1151+
Loc);
1152+
if (Res.first.isInvalid() && Res.second == OpenACCParseCanContinue::Cannot)
11481153
return true;
11491154

11501155
if (ExpectAndConsume(tok::colon))
@@ -1172,9 +1177,13 @@ bool Parser::ParseOpenACCWaitArgument(SourceLocation Loc, bool IsDirective) {
11721177
}
11731178
FirstArg = false;
11741179

1175-
ExprResult CurArg = ParseOpenACCAsyncArgument();
1180+
OpenACCIntExprParseResult Res = ParseOpenACCAsyncArgument(
1181+
IsDirective ? OpenACCDirectiveKind::Wait
1182+
: OpenACCDirectiveKind::Invalid,
1183+
IsDirective ? OpenACCClauseKind::Invalid : OpenACCClauseKind::Wait,
1184+
Loc);
11761185

1177-
if (CurArg.isInvalid())
1186+
if (Res.first.isInvalid() && Res.second == OpenACCParseCanContinue::Cannot)
11781187
return true;
11791188
}
11801189

‎clang/lib/Sema/SemaOpenACC.cpp

Lines changed: 40 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -198,6 +198,25 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
198198
default:
199199
return false;
200200
}
201+
case OpenACCClauseKind::Async:
202+
switch (DirectiveKind) {
203+
case OpenACCDirectiveKind::Parallel:
204+
case OpenACCDirectiveKind::Serial:
205+
case OpenACCDirectiveKind::Kernels:
206+
case OpenACCDirectiveKind::Data:
207+
case OpenACCDirectiveKind::EnterData:
208+
case OpenACCDirectiveKind::ExitData:
209+
case OpenACCDirectiveKind::Set:
210+
case OpenACCDirectiveKind::Update:
211+
case OpenACCDirectiveKind::Wait:
212+
case OpenACCDirectiveKind::ParallelLoop:
213+
case OpenACCDirectiveKind::SerialLoop:
214+
case OpenACCDirectiveKind::KernelsLoop:
215+
return true;
216+
default:
217+
return false;
218+
}
219+
201220
default:
202221
// Do nothing so we can go to the 'unimplemented' diagnostic instead.
203222
return true;
@@ -398,6 +417,27 @@ SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses,
398417
getASTContext(), Clause.getBeginLoc(), Clause.getLParenLoc(),
399418
Clause.getIntExprs()[0], Clause.getEndLoc());
400419
}
420+
case OpenACCClauseKind::Async: {
421+
// Restrictions only properly implemented on 'compute' constructs, and
422+
// 'compute' constructs are the only construct that can do anything with
423+
// this yet, so skip/treat as unimplemented in this case.
424+
if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()))
425+
break;
426+
427+
// There is no prose in the standard that says duplicates aren't allowed,
428+
// but this diagnostic is present in other compilers, as well as makes
429+
// sense.
430+
if (checkAlreadyHasClauseOfKind(*this, ExistingClauses, Clause))
431+
return nullptr;
432+
433+
assert(Clause.getNumIntExprs() < 2 &&
434+
"Invalid number of expressions for Async");
435+
436+
return OpenACCAsyncClause::Create(
437+
getASTContext(), Clause.getBeginLoc(), Clause.getLParenLoc(),
438+
Clause.getNumIntExprs() != 0 ? Clause.getIntExprs()[0] : nullptr,
439+
Clause.getEndLoc());
440+
}
401441
case OpenACCClauseKind::Private: {
402442
// Restrictions only properly implemented on 'compute' constructs, and
403443
// 'compute' constructs are the only construct that can do anything with

‎clang/lib/Sema/TreeTransform.h

Lines changed: 24 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -11401,6 +11401,30 @@ void OpenACCClauseTransform<Derived>::VisitVectorLengthClause(
1140111401
ParsedClause.getLParenLoc(), ParsedClause.getIntExprs()[0],
1140211402
ParsedClause.getEndLoc());
1140311403
}
11404+
11405+
template <typename Derived>
11406+
void OpenACCClauseTransform<Derived>::VisitAsyncClause(
11407+
const OpenACCAsyncClause &C) {
11408+
if (C.hasIntExpr()) {
11409+
ExprResult Res = Self.TransformExpr(const_cast<Expr *>(C.getIntExpr()));
11410+
if (!Res.isUsable())
11411+
return;
11412+
11413+
Res = Self.getSema().OpenACC().ActOnIntExpr(OpenACCDirectiveKind::Invalid,
11414+
C.getClauseKind(),
11415+
C.getBeginLoc(), Res.get());
11416+
if (!Res.isUsable())
11417+
return;
11418+
ParsedClause.setIntExprDetails(Res.get());
11419+
}
11420+
11421+
NewClause = OpenACCAsyncClause::Create(
11422+
Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
11423+
ParsedClause.getLParenLoc(),
11424+
ParsedClause.getNumIntExprs() != 0 ? ParsedClause.getIntExprs()[0]
11425+
: nullptr,
11426+
ParsedClause.getEndLoc());
11427+
}
1140411428
} // namespace
1140511429
template <typename Derived>
1140611430
OpenACCClause *TreeTransform<Derived>::TransformOpenACCClause(

‎clang/lib/Serialization/ASTReader.cpp

Lines changed: 6 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -11881,6 +11881,12 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
1188111881
return OpenACCCreateClause::Create(getContext(), ClauseKind, BeginLoc,
1188211882
LParenLoc, IsZero, VarList, EndLoc);
1188311883
}
11884+
case OpenACCClauseKind::Async: {
11885+
SourceLocation LParenLoc = readSourceLocation();
11886+
Expr *AsyncExpr = readBool() ? readSubExpr() : nullptr;
11887+
return OpenACCAsyncClause::Create(getContext(), BeginLoc, LParenLoc,
11888+
AsyncExpr, EndLoc);
11889+
}
1188411890

1188511891
case OpenACCClauseKind::Finalize:
1188611892
case OpenACCClauseKind::IfPresent:
@@ -11904,7 +11910,6 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
1190411910
case OpenACCClauseKind::DefaultAsync:
1190511911
case OpenACCClauseKind::DeviceType:
1190611912
case OpenACCClauseKind::DType:
11907-
case OpenACCClauseKind::Async:
1190811913
case OpenACCClauseKind::Tile:
1190911914
case OpenACCClauseKind::Gang:
1191011915
case OpenACCClauseKind::Wait:

‎clang/lib/Serialization/ASTWriter.cpp

Lines changed: 8 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -7908,6 +7908,14 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
79087908
writeOpenACCVarList(CC);
79097909
return;
79107910
}
7911+
case OpenACCClauseKind::Async: {
7912+
const auto *AC = cast<OpenACCAsyncClause>(C);
7913+
writeSourceLocation(AC->getLParenLoc());
7914+
writeBool(AC->hasIntExpr());
7915+
if (AC->hasIntExpr())
7916+
AddStmt(const_cast<Expr *>(AC->getIntExpr()));
7917+
return;
7918+
}
79117919

79127920
case OpenACCClauseKind::Finalize:
79137921
case OpenACCClauseKind::IfPresent:
@@ -7931,7 +7939,6 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
79317939
case OpenACCClauseKind::DefaultAsync:
79327940
case OpenACCClauseKind::DeviceType:
79337941
case OpenACCClauseKind::DType:
7934-
case OpenACCClauseKind::Async:
79357942
case OpenACCClauseKind::Tile:
79367943
case OpenACCClauseKind::Gang:
79377944
case OpenACCClauseKind::Wait:

‎clang/test/AST/ast-print-openacc-compute-construct.cpp

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -75,5 +75,13 @@ void foo() {
7575
// CHECK: #pragma acc kernels deviceptr(iPtr, arrayPtr[0])
7676
#pragma acc kernels deviceptr(iPtr, arrayPtr[0])
7777
while(true);
78+
79+
// CHECK: #pragma acc kernels async(*iPtr)
80+
#pragma acc kernels async(*iPtr)
81+
while(true);
82+
83+
// CHECK: #pragma acc kernels async
84+
#pragma acc kernels async
85+
while(true);
7886
}
7987

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

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1233,7 +1233,6 @@ void device_type() {
12331233

12341234
#define acc_async_sync -1
12351235
void AsyncArgument() {
1236-
// expected-warning@+1{{OpenACC clause 'async' not yet implemented, clause ignored}}
12371236
#pragma acc parallel async
12381237
{}
12391238

@@ -1250,15 +1249,12 @@ void AsyncArgument() {
12501249
#pragma acc parallel async(4, 3)
12511250
{}
12521251

1253-
// expected-warning@+1{{OpenACC clause 'async' not yet implemented, clause ignored}}
12541252
#pragma acc parallel async(returns_int())
12551253
{}
12561254

1257-
// expected-warning@+1{{OpenACC clause 'async' not yet implemented, clause ignored}}
12581255
#pragma acc parallel async(5)
12591256
{}
12601257

1261-
// expected-warning@+1{{OpenACC clause 'async' not yet implemented, clause ignored}}
12621258
#pragma acc parallel async(acc_async_sync)
12631259
{}
12641260
}

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

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -18,13 +18,14 @@ void templ() {
1818
#pragma acc parallel vector_length(I)
1919
for(;;){}
2020

21-
// expected-warning@+1{{OpenACC clause 'async' not yet implemented, clause ignored}}
2221
#pragma acc parallel async(T::value)
2322
for(;;){}
2423

25-
// expected-warning@+1{{OpenACC clause 'async' not yet implemented, clause ignored}}
2624
#pragma acc parallel async(I)
2725
for(;;){}
26+
27+
#pragma acc parallel async
28+
for(;;){}
2829
}
2930

3031
struct S {

0 commit comments

Comments
 (0)