Skip to content

Commit 21c785d

Browse files
erichkeaneerichkeane
committed
[OpenACC] Implement 'set' construct sema
The 'set' construct is another fairly simple one, it doesn't have an associated statement and only a handful of allowed clauses. This patch implements it and all the rules for it, allowing 3 of its for clauses. The only exception is default_async, which will be implemented in a future patch, because it isn't just being enabled, it needs a complete new implementation.
1 parent 3f93625 commit 21c785d

26 files changed

+384
-16
lines changed

‎clang/include/clang-c/Index.h

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2198,7 +2198,11 @@ enum CXCursorKind {
21982198
*/
21992199
CXCursor_OpenACCShutdownConstruct = 329,
22002200

2201-
CXCursor_LastStmt = CXCursor_OpenACCShutdownConstruct,
2201+
/** OpenACC set Construct.
2202+
*/
2203+
CXCursor_OpenACCSetConstruct = 330,
2204+
2205+
CXCursor_LastStmt = CXCursor_OpenACCSetConstruct,
22022206

22032207
/**
22042208
* Cursor that represents the translation unit itself.

‎clang/include/clang/AST/RecursiveASTVisitor.h

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4080,6 +4080,8 @@ DEF_TRAVERSE_STMT(OpenACCInitConstruct,
40804080
{ TRY_TO(VisitOpenACCClauseList(S->clauses())); })
40814081
DEF_TRAVERSE_STMT(OpenACCShutdownConstruct,
40824082
{ TRY_TO(VisitOpenACCClauseList(S->clauses())); })
4083+
DEF_TRAVERSE_STMT(OpenACCSetConstruct,
4084+
{ TRY_TO(VisitOpenACCClauseList(S->clauses())); })
40834085

40844086
// Traverse HLSL: Out argument expression
40854087
DEF_TRAVERSE_STMT(HLSLOutArgExpr, {})

‎clang/include/clang/AST/StmtOpenACC.h

Lines changed: 40 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -672,5 +672,45 @@ class OpenACCShutdownConstruct final
672672
SourceLocation End, ArrayRef<const OpenACCClause *> Clauses);
673673
};
674674

675+
// This class represents a 'set' construct, which has just a clause list.
676+
class OpenACCSetConstruct final
677+
: public OpenACCConstructStmt,
678+
private llvm::TrailingObjects<OpenACCSetConstruct,
679+
const OpenACCClause *> {
680+
friend TrailingObjects;
681+
OpenACCSetConstruct(unsigned NumClauses)
682+
: OpenACCConstructStmt(OpenACCSetConstructClass,
683+
OpenACCDirectiveKind::Set, SourceLocation{},
684+
SourceLocation{}, SourceLocation{}) {
685+
std::uninitialized_value_construct(
686+
getTrailingObjects<const OpenACCClause *>(),
687+
getTrailingObjects<const OpenACCClause *>() + NumClauses);
688+
setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
689+
NumClauses));
690+
}
691+
692+
OpenACCSetConstruct(SourceLocation Start, SourceLocation DirectiveLoc,
693+
SourceLocation End,
694+
ArrayRef<const OpenACCClause *> Clauses)
695+
: OpenACCConstructStmt(OpenACCSetConstructClass,
696+
OpenACCDirectiveKind::Set, Start, DirectiveLoc,
697+
End) {
698+
std::uninitialized_copy(Clauses.begin(), Clauses.end(),
699+
getTrailingObjects<const OpenACCClause *>());
700+
setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
701+
Clauses.size()));
702+
}
703+
704+
public:
705+
static bool classof(const Stmt *T) {
706+
return T->getStmtClass() == OpenACCSetConstructClass;
707+
}
708+
static OpenACCSetConstruct *CreateEmpty(const ASTContext &C,
709+
unsigned NumClauses);
710+
static OpenACCSetConstruct *Create(const ASTContext &C, SourceLocation Start,
711+
SourceLocation DirectiveLoc,
712+
SourceLocation End,
713+
ArrayRef<const OpenACCClause *> Clauses);
714+
};
675715
} // namespace clang
676716
#endif // LLVM_CLANG_AST_STMTOPENACC_H

‎clang/include/clang/AST/TextNodeDumper.h

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -417,6 +417,7 @@ class TextNodeDumper
417417
void VisitOpenACCHostDataConstruct(const OpenACCHostDataConstruct *S);
418418
void VisitOpenACCWaitConstruct(const OpenACCWaitConstruct *S);
419419
void VisitOpenACCInitConstruct(const OpenACCInitConstruct *S);
420+
void VisitOpenACCSetConstruct(const OpenACCSetConstruct *S);
420421
void VisitOpenACCShutdownConstruct(const OpenACCShutdownConstruct *S);
421422
void VisitOpenACCAsteriskSizeExpr(const OpenACCAsteriskSizeExpr *S);
422423
void VisitEmbedExpr(const EmbedExpr *S);

‎clang/include/clang/Basic/StmtNodes.td

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -315,6 +315,7 @@ def OpenACCHostDataConstruct : StmtNode<OpenACCAssociatedStmtConstruct>;
315315
def OpenACCWaitConstruct : StmtNode<OpenACCConstructStmt>;
316316
def OpenACCInitConstruct : StmtNode<OpenACCConstructStmt>;
317317
def OpenACCShutdownConstruct : StmtNode<OpenACCConstructStmt>;
318+
def OpenACCSetConstruct : StmtNode<OpenACCConstructStmt>;
318319

319320
// OpenACC Additional Expressions.
320321
def OpenACCAsteriskSizeExpr : StmtNode<Expr>;

‎clang/include/clang/Serialization/ASTBitCodes.h

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2024,6 +2024,7 @@ enum StmtCode {
20242024
STMT_OPENACC_WAIT_CONSTRUCT,
20252025
STMT_OPENACC_INIT_CONSTRUCT,
20262026
STMT_OPENACC_SHUTDOWN_CONSTRUCT,
2027+
STMT_OPENACC_SET_CONSTRUCT,
20272028

20282029
// HLSL Constructs
20292030
EXPR_HLSL_OUT_ARG,

‎clang/lib/AST/StmtOpenACC.cpp

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -265,3 +265,22 @@ OpenACCShutdownConstruct *OpenACCShutdownConstruct::Create(
265265
new (Mem) OpenACCShutdownConstruct(Start, DirectiveLoc, End, Clauses);
266266
return Inst;
267267
}
268+
269+
OpenACCSetConstruct *OpenACCSetConstruct::CreateEmpty(const ASTContext &C,
270+
unsigned NumClauses) {
271+
void *Mem = C.Allocate(
272+
OpenACCSetConstruct::totalSizeToAlloc<const OpenACCClause *>(NumClauses));
273+
auto *Inst = new (Mem) OpenACCSetConstruct(NumClauses);
274+
return Inst;
275+
}
276+
277+
OpenACCSetConstruct *
278+
OpenACCSetConstruct::Create(const ASTContext &C, SourceLocation Start,
279+
SourceLocation DirectiveLoc, SourceLocation End,
280+
ArrayRef<const OpenACCClause *> Clauses) {
281+
void *Mem =
282+
C.Allocate(OpenACCSetConstruct::totalSizeToAlloc<const OpenACCClause *>(
283+
Clauses.size()));
284+
auto *Inst = new (Mem) OpenACCSetConstruct(Start, DirectiveLoc, End, Clauses);
285+
return Inst;
286+
}

‎clang/lib/AST/StmtPrinter.cpp

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1205,6 +1205,10 @@ void StmtPrinter::VisitOpenACCShutdownConstruct(OpenACCShutdownConstruct *S) {
12051205
PrintOpenACCConstruct(S);
12061206
}
12071207

1208+
void StmtPrinter::VisitOpenACCSetConstruct(OpenACCSetConstruct *S) {
1209+
PrintOpenACCConstruct(S);
1210+
}
1211+
12081212
void StmtPrinter::VisitOpenACCWaitConstruct(OpenACCWaitConstruct *S) {
12091213
Indent() << "#pragma acc wait";
12101214
if (!S->getLParenLoc().isInvalid()) {

‎clang/lib/AST/StmtProfile.cpp

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2769,6 +2769,12 @@ void StmtProfiler::VisitOpenACCShutdownConstruct(
27692769
P.VisitOpenACCClauseList(S->clauses());
27702770
}
27712771

2772+
void StmtProfiler::VisitOpenACCSetConstruct(const OpenACCSetConstruct *S) {
2773+
VisitStmt(S);
2774+
OpenACCClauseProfiler P{*this};
2775+
P.VisitOpenACCClauseList(S->clauses());
2776+
}
2777+
27722778
void StmtProfiler::VisitHLSLOutArgExpr(const HLSLOutArgExpr *S) {
27732779
VisitStmt(S);
27742780
}

‎clang/lib/AST/TextNodeDumper.cpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2971,6 +2971,9 @@ void TextNodeDumper::VisitOpenACCShutdownConstruct(
29712971
const OpenACCShutdownConstruct *S) {
29722972
OS << " " << S->getDirectiveKind();
29732973
}
2974+
void TextNodeDumper::VisitOpenACCSetConstruct(const OpenACCSetConstruct *S) {
2975+
OS << " " << S->getDirectiveKind();
2976+
}
29742977

29752978
void TextNodeDumper::VisitEmbedExpr(const EmbedExpr *S) {
29762979
AddChild("begin", [=] { OS << S->getStartingElementPos(); });

‎clang/lib/CodeGen/CGStmt.cpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -479,6 +479,9 @@ void CodeGenFunction::EmitStmt(const Stmt *S, ArrayRef<const Attr *> Attrs) {
479479
case Stmt::OpenACCShutdownConstructClass:
480480
EmitOpenACCShutdownConstruct(cast<OpenACCShutdownConstruct>(*S));
481481
break;
482+
case Stmt::OpenACCSetConstructClass:
483+
EmitOpenACCSetConstruct(cast<OpenACCSetConstruct>(*S));
484+
break;
482485
}
483486
}
484487

‎clang/lib/CodeGen/CodeGenFunction.h

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4137,6 +4137,11 @@ class CodeGenFunction : public CodeGenTypeCache {
41374137
// but in the future we will implement some sort of IR.
41384138
}
41394139

4140+
void EmitOpenACCSetConstruct(const OpenACCSetConstruct &S) {
4141+
// TODO OpenACC: Implement this. It is currently implemented as a 'no-op',
4142+
// but in the future we will implement some sort of IR.
4143+
}
4144+
41404145
//===--------------------------------------------------------------------===//
41414146
// LValue Expression Emission
41424147
//===--------------------------------------------------------------------===//

‎clang/lib/Sema/SemaExceptionSpec.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1401,6 +1401,7 @@ CanThrowResult Sema::canThrow(const Stmt *S) {
14011401
case Stmt::OpenACCWaitConstructClass:
14021402
case Stmt::OpenACCInitConstructClass:
14031403
case Stmt::OpenACCShutdownConstructClass:
1404+
case Stmt::OpenACCSetConstructClass:
14041405
// These expressions can never throw.
14051406
return CT_Cannot;
14061407

‎clang/lib/Sema/SemaOpenACC.cpp

Lines changed: 42 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -587,7 +587,8 @@ bool isDirectiveKindImplemented(OpenACCDirectiveKind DK) {
587587
isOpenACCCombinedDirectiveKind(DK) || isOpenACCDataDirectiveKind(DK) ||
588588
DK == OpenACCDirectiveKind::Loop || DK == OpenACCDirectiveKind::Wait ||
589589
DK == OpenACCDirectiveKind::Init ||
590-
DK == OpenACCDirectiveKind::Shutdown;
590+
DK == OpenACCDirectiveKind::Shutdown ||
591+
DK == OpenACCDirectiveKind::Set;
591592
}
592593

593594
class SemaOpenACCClauseVisitor {
@@ -709,9 +710,9 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitIfClause(
709710

710711
// There is no prose in the standard that says duplicates aren't allowed,
711712
// but this diagnostic is present in other compilers, as well as makes
712-
// sense. Prose DOES exist for 'data' and 'host_data', 'enter data' and 'exit
713-
// data' both don't, but other implmementations do this. OpenACC issue 519
714-
// filed for the latter two.
713+
// sense. Prose DOES exist for 'data' and 'host_data', 'set', 'enter data' and
714+
// 'exit data' both don't, but other implmementations do this. OpenACC issue
715+
// 519 filed for the latter two.
715716
// GCC allows this on init/shutdown, presumably for good reason, so we do too.
716717
if (Clause.getDirectiveKind() != OpenACCDirectiveKind::Init &&
717718
Clause.getDirectiveKind() != OpenACCDirectiveKind::Shutdown &&
@@ -963,6 +964,12 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitDeviceNumClause(
963964
if (!isDirectiveKindImplemented(Clause.getDirectiveKind()))
964965
return isNotImplemented();
965966

967+
// OpenACC 3.3 2.14.3: Two instances of the same clause may not appear on the
968+
// same directive.
969+
if (Clause.getDirectiveKind() == OpenACCDirectiveKind::Set &&
970+
checkAlreadyHasClauseOfKind(SemaRef, ExistingClauses, Clause))
971+
return nullptr;
972+
966973
assert(Clause.getNumIntExprs() == 1 &&
967974
"Invalid number of expressions for device_num");
968975
return OpenACCDeviceNumClause::Create(
@@ -1177,6 +1184,12 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitDeviceTypeClause(
11771184
if (!isDirectiveKindImplemented(Clause.getDirectiveKind()))
11781185
return isNotImplemented();
11791186

1187+
// OpenACC 3.3 2.14.3: Two instances of the same clause may not appear on the
1188+
// same directive.
1189+
if (Clause.getDirectiveKind() == OpenACCDirectiveKind::Set &&
1190+
checkAlreadyHasClauseOfKind(SemaRef, ExistingClauses, Clause))
1191+
return nullptr;
1192+
11801193
// TODO OpenACC: Once we get enough of the CodeGen implemented that we have
11811194
// a source for the list of valid architectures, we need to warn on unknown
11821195
// identifiers here.
@@ -1900,6 +1913,7 @@ bool PreserveLoopRAIIDepthInAssociatedStmtRAII(OpenACCDirectiveKind DK) {
19001913
case OpenACCDirectiveKind::Wait:
19011914
case OpenACCDirectiveKind::Init:
19021915
case OpenACCDirectiveKind::Shutdown:
1916+
case OpenACCDirectiveKind::Set:
19031917
llvm_unreachable("Doesn't have an associated stmt");
19041918
default:
19051919
case OpenACCDirectiveKind::Invalid:
@@ -2328,6 +2342,7 @@ void SemaOpenACC::ActOnConstruct(OpenACCDirectiveKind K,
23282342
case OpenACCDirectiveKind::HostData:
23292343
case OpenACCDirectiveKind::Init:
23302344
case OpenACCDirectiveKind::Shutdown:
2345+
case OpenACCDirectiveKind::Set:
23312346
// Nothing to do here, there is no real legalization that needs to happen
23322347
// here as these constructs do not take any arguments.
23332348
break;
@@ -3661,6 +3676,24 @@ bool SemaOpenACC::ActOnStartStmtDirective(
36613676
return Diag(StartLoc, diag::err_acc_construct_one_clause_of)
36623677
<< K << GetListOfClauses({OpenACCClauseKind::UseDevice});
36633678

3679+
// OpenACC3.3 2.14.3: At least one default_async, device_num, or device_type
3680+
// clause must appear.
3681+
if (K == OpenACCDirectiveKind::Set &&
3682+
llvm::find_if(
3683+
Clauses,
3684+
llvm::IsaPred</*OpenACCDefaultAsyncClause,*/ // TODO: ERICH Need to
3685+
// implement.
3686+
// DefaultAsyncClause
3687+
// then enable this.
3688+
OpenACCDeviceNumClause, OpenACCDeviceTypeClause,
3689+
OpenACCIfClause>) == Clauses.end())
3690+
return Diag(StartLoc, diag::err_acc_construct_one_clause_of)
3691+
<< K
3692+
<< GetListOfClauses({OpenACCClauseKind::DefaultAsync,
3693+
OpenACCClauseKind::DeviceNum,
3694+
OpenACCClauseKind::DeviceType,
3695+
OpenACCClauseKind::If});
3696+
36643697
return diagnoseConstructAppertainment(*this, K, StartLoc, /*IsStmt=*/true);
36653698
}
36663699

@@ -3724,6 +3757,10 @@ StmtResult SemaOpenACC::ActOnEndStmtDirective(
37243757
return OpenACCShutdownConstruct::Create(getASTContext(), StartLoc, DirLoc,
37253758
EndLoc, Clauses);
37263759
}
3760+
case OpenACCDirectiveKind::Set: {
3761+
return OpenACCSetConstruct::Create(getASTContext(), StartLoc, DirLoc,
3762+
EndLoc, Clauses);
3763+
}
37273764
}
37283765
llvm_unreachable("Unhandled case in directive handling?");
37293766
}
@@ -3739,6 +3776,7 @@ StmtResult SemaOpenACC::ActOnAssociatedStmt(
37393776
case OpenACCDirectiveKind::Wait:
37403777
case OpenACCDirectiveKind::Init:
37413778
case OpenACCDirectiveKind::Shutdown:
3779+
case OpenACCDirectiveKind::Set:
37423780
llvm_unreachable(
37433781
"these don't have associated statements, so shouldn't get here");
37443782
case OpenACCDirectiveKind::Parallel:

‎clang/lib/Sema/TreeTransform.h

Lines changed: 25 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4169,6 +4169,15 @@ class TreeTransform {
41694169
SourceLocation{}, {}, SourceLocation{}, EndLoc, Clauses, {});
41704170
}
41714171

4172+
StmtResult RebuildOpenACCSetConstruct(SourceLocation BeginLoc,
4173+
SourceLocation DirLoc,
4174+
SourceLocation EndLoc,
4175+
ArrayRef<OpenACCClause *> Clauses) {
4176+
return getSema().OpenACC().ActOnEndStmtDirective(
4177+
OpenACCDirectiveKind::Set, BeginLoc, DirLoc, SourceLocation{},
4178+
SourceLocation{}, {}, SourceLocation{}, EndLoc, Clauses, {});
4179+
}
4180+
41724181
StmtResult RebuildOpenACCWaitConstruct(
41734182
SourceLocation BeginLoc, SourceLocation DirLoc, SourceLocation LParenLoc,
41744183
Expr *DevNumExpr, SourceLocation QueuesLoc, ArrayRef<Expr *> QueueIdExprs,
@@ -12422,6 +12431,22 @@ StmtResult TreeTransform<Derived>::TransformOpenACCShutdownConstruct(
1242212431
C->getBeginLoc(), C->getDirectiveLoc(), C->getEndLoc(),
1242312432
TransformedClauses);
1242412433
}
12434+
template <typename Derived>
12435+
StmtResult
12436+
TreeTransform<Derived>::TransformOpenACCSetConstruct(OpenACCSetConstruct *C) {
12437+
getSema().OpenACC().ActOnConstruct(C->getDirectiveKind(), C->getBeginLoc());
12438+
12439+
llvm::SmallVector<OpenACCClause *> TransformedClauses =
12440+
getDerived().TransformOpenACCClauseList(C->getDirectiveKind(),
12441+
C->clauses());
12442+
if (getSema().OpenACC().ActOnStartStmtDirective(
12443+
C->getDirectiveKind(), C->getBeginLoc(), TransformedClauses))
12444+
return StmtError();
12445+
12446+
return getDerived().RebuildOpenACCSetConstruct(
12447+
C->getBeginLoc(), C->getDirectiveLoc(), C->getEndLoc(),
12448+
TransformedClauses);
12449+
}
1242512450

1242612451
template <typename Derived>
1242712452
StmtResult

‎clang/lib/Serialization/ASTReaderStmt.cpp

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2875,6 +2875,11 @@ void ASTStmtReader::VisitOpenACCShutdownConstruct(OpenACCShutdownConstruct *S) {
28752875
VisitOpenACCConstructStmt(S);
28762876
}
28772877

2878+
void ASTStmtReader::VisitOpenACCSetConstruct(OpenACCSetConstruct *S) {
2879+
VisitStmt(S);
2880+
VisitOpenACCConstructStmt(S);
2881+
}
2882+
28782883
void ASTStmtReader::VisitOpenACCHostDataConstruct(OpenACCHostDataConstruct *S) {
28792884
VisitStmt(S);
28802885
VisitOpenACCAssociatedStmtConstruct(S);
@@ -4407,6 +4412,11 @@ Stmt *ASTReader::ReadStmtFromStream(ModuleFile &F) {
44074412
S = OpenACCShutdownConstruct::CreateEmpty(Context, NumClauses);
44084413
break;
44094414
}
4415+
case STMT_OPENACC_SET_CONSTRUCT: {
4416+
unsigned NumClauses = Record[ASTStmtReader::NumStmtFields];
4417+
S = OpenACCSetConstruct::CreateEmpty(Context, NumClauses);
4418+
break;
4419+
}
44104420
case EXPR_REQUIRES: {
44114421
unsigned numLocalParameters = Record[ASTStmtReader::NumExprFields];
44124422
unsigned numRequirement = Record[ASTStmtReader::NumExprFields + 1];

‎clang/lib/Serialization/ASTWriterStmt.cpp

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2957,6 +2957,12 @@ void ASTStmtWriter::VisitOpenACCShutdownConstruct(OpenACCShutdownConstruct *S) {
29572957
Code = serialization::STMT_OPENACC_SHUTDOWN_CONSTRUCT;
29582958
}
29592959

2960+
void ASTStmtWriter::VisitOpenACCSetConstruct(OpenACCSetConstruct *S) {
2961+
VisitStmt(S);
2962+
VisitOpenACCConstructStmt(S);
2963+
Code = serialization::STMT_OPENACC_SET_CONSTRUCT;
2964+
}
2965+
29602966
void ASTStmtWriter::VisitOpenACCHostDataConstruct(OpenACCHostDataConstruct *S) {
29612967
VisitStmt(S);
29622968
VisitOpenACCAssociatedStmtConstruct(S);

‎clang/lib/StaticAnalyzer/Core/ExprEngine.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1832,6 +1832,7 @@ void ExprEngine::Visit(const Stmt *S, ExplodedNode *Pred,
18321832
case Stmt::OpenACCWaitConstructClass:
18331833
case Stmt::OpenACCInitConstructClass:
18341834
case Stmt::OpenACCShutdownConstructClass:
1835+
case Stmt::OpenACCSetConstructClass:
18351836
case Stmt::OMPUnrollDirectiveClass:
18361837
case Stmt::OMPMetaDirectiveClass:
18371838
case Stmt::HLSLOutArgExprClass: {

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

Lines changed: 25 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,25 @@
1+
// RUN: %clang_cc1 -fopenacc -ast-print %s -o - | FileCheck %s
2+
3+
unsigned Int;
4+
5+
void uses() {
6+
// CHECK: #pragma acc set if(Int == 5) device_type(I) device_num(Int)
7+
#pragma acc set default_async(Int) if (Int == 5) device_type(I) device_num(Int)
8+
// CHECK: #pragma acc set device_type(I) device_num(Int)
9+
#pragma acc set default_async(Int) device_type(I) device_num(Int)
10+
// CHECK: #pragma acc set if(Int == 5) device_num(Int)
11+
#pragma acc set default_async(Int) if (Int == 5) device_num(Int)
12+
// CHECK: #pragma acc set if(Int == 5) device_type(I)
13+
#pragma acc set default_async(Int) if (Int == 5) device_type(I)
14+
// CHECK: #pragma acc set if(Int == 5) device_type(I) device_num(Int)
15+
#pragma acc set if (Int == 5) device_type(I) device_num(Int)
16+
// CHECK-NOT: default_async(Int)
17+
// TODO: OpenACC: Enable this when we have default_async implemented.
18+
// #pragma acc set default_async(Int)
19+
// CHECK: #pragma acc set if(Int == 5)
20+
#pragma acc set if (Int == 5)
21+
// CHECK: #pragma acc set device_type(I)
22+
#pragma acc set device_type(I)
23+
// CHECK: #pragma acc set device_num(Int)
24+
#pragma acc set device_num(Int)
25+
}

0 commit comments

Comments
 (0)