diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index c29a3422acd26..90cf1bd9b2588 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12961,6 +12961,19 @@ def err_acc_clause_after_device_type def err_acc_clause_cannot_combine : Error<"OpenACC clause '%0' may not appear on the same construct as a " "'%1' clause on a '%2' construct">; +def err_acc_clause_routine_cannot_combine_before_device_type + : Error<"OpenACC clause '%0' after 'device_type' clause on a 'routine' " + "conflicts with the '%1' clause before the first 'device_type'">; +def err_acc_clause_routine_cannot_combine_same_device_type + : Error<"OpenACC clause '%0' on a 'routine' directive conflicts with the " + "'%1' clause applying to the same 'device_type'">; +def err_acc_clause_routine_one_of_in_region + : Error<"OpenACC 'routine' construct must have at least one 'gang', 'seq', " + "'vector', or 'worker' clause that applies to each 'device_type'">; +def err_acc_clause_since_last_device_type + : Error<"OpenACC '%0' clause cannot appear more than once%select{| in a " + "'device_type' region}2 on a '%1' directive">; + def err_acc_reduction_num_gangs_conflict : Error<"OpenACC '%1' clause %select{|with more than 1 argument }0may not " "appear on a '%2' construct " @@ -12978,9 +12991,6 @@ def note_acc_reduction_composite_member_loc : Note<"invalid field is here">; def err_acc_loop_not_for_loop : Error<"OpenACC '%0' construct can only be applied to a 'for' loop">; def note_acc_construct_here : Note<"'%0' construct is here">; -def err_acc_loop_spec_conflict - : Error<"OpenACC clause '%0' on '%1' construct conflicts with previous " - "data dependence clause">; def err_acc_collapse_loop_count : Error<"OpenACC 'collapse' clause loop count must be a %select{constant " "expression|positive integer value, evaluated to %1}0">; @@ -13032,9 +13042,10 @@ def err_acc_gang_reduction_numgangs_conflict "'%1' clause %select{inside a compute construct with a|and a}3 " "'num_gangs' clause with more than one argument">; - def err_reduction_op_mismatch +def err_reduction_op_mismatch : Error<"OpenACC 'reduction' variable must have the same operator in all " "nested constructs (%0 vs %1)">; + def err_acc_loop_variable_type : Error<"loop variable of loop associated with an OpenACC '%0' construct " "must be of integer, pointer, or random-access-iterator type (is " diff --git a/clang/include/clang/Basic/OpenACCKinds.h b/clang/include/clang/Basic/OpenACCKinds.h index 652831e23a758..a798401dfa9f5 100644 --- a/clang/include/clang/Basic/OpenACCKinds.h +++ b/clang/include/clang/Basic/OpenACCKinds.h @@ -202,32 +202,21 @@ inline llvm::raw_ostream &operator<<(llvm::raw_ostream &Out, return printOpenACCAtomicKind(Out, AK); } -/// Represents the kind of an OpenACC clause. +/// Represents the kind of an OpenACC clause. Sorted alphabetically, since this +/// order ends up influencing the sorting of the list diagnostic. enum class OpenACCClauseKind : uint8_t { - /// 'finalize' clause, allowed on 'exit data' directive. - Finalize, - /// 'if_present' clause, allowed on 'host_data' and 'update' directives. - IfPresent, - /// 'seq' clause, allowed on 'loop' and 'routine' directives. - Seq, - /// 'independent' clause, allowed on 'loop' directives. - Independent, + /// 'async' clause, allowed on Compute, Data, 'update', 'wait', and Combined + /// constructs. + Async, + /// 'attach' clause, allowed on Compute and Combined constructs, plus 'data' + /// and 'enter data'. + Attach, /// 'auto' clause, allowed on 'loop' directives. Auto, - /// 'worker' clause, allowed on 'loop', Combined, and 'routine' directives. - Worker, - /// 'vector' clause, allowed on 'loop', Combined, and 'routine' directives. - Vector, - /// 'nohost' clause, allowed on 'routine' directives. - NoHost, - /// 'default' clause, allowed on parallel, serial, kernel (and compound) - /// constructs. - Default, - /// 'if' clause, allowed on all the Compute Constructs, Data Constructs, - /// Executable Constructs, and Combined Constructs. - If, - /// 'self' clause, allowed on Compute and Combined Constructs, plus 'update'. - Self, + /// 'bind' clause, allowed on routine constructs. + Bind, + /// 'collapse' clause, allowed on 'loop' and Combined constructs. + Collapse, /// 'copy' clause, allowed on Compute and Combined Constructs, plus 'data' and /// 'declare'. Copy, @@ -235,100 +224,118 @@ enum class OpenACCClauseKind : uint8_t { PCopy, /// 'copy' clause alias 'present_or_copy'. Preserved for diagnostic purposes. PresentOrCopy, - /// 'use_device' clause, allowed on 'host_data' construct. - UseDevice, - /// 'attach' clause, allowed on Compute and Combined constructs, plus 'data' - /// and 'enter data'. - Attach, + /// 'copyin' clause, allowed on Compute and Combined constructs, plus 'data', + /// 'enter data', and 'declare'. + CopyIn, + /// 'copyin' clause alias 'pcopyin'. Preserved for diagnostic purposes. + PCopyIn, + /// 'copyin' clause alias 'present_or_copyin'. Preserved for diagnostic + /// purposes. + PresentOrCopyIn, + /// 'copyout' clause, allowed on Compute and Combined constructs, plus 'data', + /// 'exit data', and 'declare'. + CopyOut, + /// 'copyout' clause alias 'pcopyout'. Preserved for diagnostic purposes. + PCopyOut, + /// 'copyout' clause alias 'present_or_copyout'. Preserved for diagnostic + /// purposes. + PresentOrCopyOut, + /// 'create' clause, allowed on Compute and Combined constructs, plus 'data', + /// 'enter data', and 'declare'. + Create, + /// 'create' clause alias 'pcreate'. Preserved for diagnostic purposes. + PCreate, + /// 'create' clause alias 'present_or_create'. Preserved for diagnostic + /// purposes. + PresentOrCreate, + /// 'default' clause, allowed on parallel, serial, kernel (and compound) + /// constructs. + Default, + /// 'default_async' clause, allowed on 'set' construct. + DefaultAsync, /// 'delete' clause, allowed on the 'exit data' construct. Delete, /// 'detach' clause, allowed on the 'exit data' construct. Detach, /// 'device' clause, allowed on the 'update' construct. Device, + /// 'device_num' clause, allowed on 'init', 'shutdown', and 'set' constructs. + DeviceNum, /// 'deviceptr' clause, allowed on Compute and Combined Constructs, plus /// 'data' and 'declare'. DevicePtr, /// 'device_resident' clause, allowed on the 'declare' construct. DeviceResident, + /// 'device_type' clause, allowed on Compute, 'data', 'init', 'shutdown', + /// 'set', update', 'loop', 'routine', and Combined constructs. + DeviceType, + /// 'dtype' clause, an alias for 'device_type', stored separately for + /// diagnostic purposes. + DType, + /// 'finalize' clause, allowed on 'exit data' directive. + Finalize, /// 'firstprivate' clause, allowed on 'parallel', 'serial', 'parallel loop', /// and 'serial loop' constructs. FirstPrivate, + /// 'gang' clause, allowed on 'loop' and Combined constructs. + Gang, /// 'host' clause, allowed on 'update' construct. Host, + /// 'if' clause, allowed on all the Compute Constructs, Data Constructs, + /// Executable Constructs, and Combined Constructs. + If, + /// 'if_present' clause, allowed on 'host_data' and 'update' directives. + IfPresent, + /// 'independent' clause, allowed on 'loop' directives. + Independent, /// 'link' clause, allowed on 'declare' construct. Link, /// 'no_create' clause, allowed on allowed on Compute and Combined constructs, /// plus 'data'. NoCreate, + /// 'nohost' clause, allowed on 'routine' directives. + NoHost, + /// 'num_gangs' clause, allowed on 'parallel', 'kernels', parallel loop', and + /// 'kernels loop' constructs. + NumGangs, + /// 'num_workers' clause, allowed on 'parallel', 'kernels', parallel loop', + /// and 'kernels loop' constructs. + NumWorkers, /// 'present' clause, allowed on Compute and Combined constructs, plus 'data' /// and 'declare'. Present, /// 'private' clause, allowed on 'parallel', 'serial', 'loop', 'parallel /// loop', and 'serial loop' constructs. Private, - /// 'copyout' clause, allowed on Compute and Combined constructs, plus 'data', - /// 'exit data', and 'declare'. - CopyOut, - /// 'copyout' clause alias 'pcopyout'. Preserved for diagnostic purposes. - PCopyOut, - /// 'copyout' clause alias 'present_or_copyout'. Preserved for diagnostic - /// purposes. - PresentOrCopyOut, - /// 'copyin' clause, allowed on Compute and Combined constructs, plus 'data', - /// 'enter data', and 'declare'. - CopyIn, - /// 'copyin' clause alias 'pcopyin'. Preserved for diagnostic purposes. - PCopyIn, - /// 'copyin' clause alias 'present_or_copyin'. Preserved for diagnostic - /// purposes. - PresentOrCopyIn, - /// 'create' clause, allowed on Compute and Combined constructs, plus 'data', - /// 'enter data', and 'declare'. - Create, - /// 'create' clause alias 'pcreate'. Preserved for diagnostic purposes. - PCreate, - /// 'create' clause alias 'present_or_create'. Preserved for diagnostic - /// purposes. - PresentOrCreate, /// 'reduction' clause, allowed on Parallel, Serial, Loop, and the combined /// constructs. Reduction, - /// 'collapse' clause, allowed on 'loop' and Combined constructs. - Collapse, - /// 'bind' clause, allowed on routine constructs. - Bind, + /// 'self' clause, allowed on Compute and Combined Constructs, plus 'update'. + Self, + /// 'seq' clause, allowed on 'loop' and 'routine' directives. + Seq, + /// 'tile' clause, allowed on 'loop' and Combined constructs. + Tile, + /// 'use_device' clause, allowed on 'host_data' construct. + UseDevice, + /// 'vector' clause, allowed on 'loop', Combined, and 'routine' directives. + Vector, /// 'vector_length' clause, allowed on 'parallel', 'kernels', 'parallel loop', /// and 'kernels loop' constructs. VectorLength, - /// 'num_gangs' clause, allowed on 'parallel', 'kernels', parallel loop', and - /// 'kernels loop' constructs. - NumGangs, - /// 'num_workers' clause, allowed on 'parallel', 'kernels', parallel loop', - /// and 'kernels loop' constructs. - NumWorkers, - /// 'device_num' clause, allowed on 'init', 'shutdown', and 'set' constructs. - DeviceNum, - /// 'default_async' clause, allowed on 'set' construct. - DefaultAsync, - /// 'device_type' clause, allowed on Compute, 'data', 'init', 'shutdown', - /// 'set', update', 'loop', 'routine', and Combined constructs. - DeviceType, - /// 'dtype' clause, an alias for 'device_type', stored separately for - /// diagnostic purposes. - DType, - /// 'async' clause, allowed on Compute, Data, 'update', 'wait', and Combined - /// constructs. - Async, - /// 'tile' clause, allowed on 'loop' and Combined constructs. - Tile, - /// 'gang' clause, allowed on 'loop' and Combined constructs. - Gang, /// 'wait' clause, allowed on Compute, Data, 'update', and Combined /// constructs. Wait, + /// 'worker' clause, allowed on 'loop', Combined, and 'routine' directives. + Worker, - /// Represents an invalid clause, for the purposes of parsing. + /// 'shortloop' is represented in the ACC.td file, but isn't present in the + /// standard. This appears to be an old extension for the nvidia fortran + // compiler, but seemingly not elsewhere. Put it here as a placeholder, but it + // is never expected to be generated. + Shortloop, + /// Represents an invalid clause, for the purposes of parsing. Should be + /// 'last'. Invalid, }; @@ -485,6 +492,9 @@ inline StreamTy &printOpenACCClauseKind(StreamTy &Out, OpenACCClauseKind K) { case OpenACCClauseKind::Wait: return Out << "wait"; + case OpenACCClauseKind::Shortloop: + llvm_unreachable("Shortloop shouldn't be generated in clang"); + LLVM_FALLTHROUGH; case OpenACCClauseKind::Invalid: return Out << ""; } diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h index 8d31d46444c7e..2c9134b0e6d33 100644 --- a/clang/include/clang/Sema/SemaOpenACC.h +++ b/clang/include/clang/Sema/SemaOpenACC.h @@ -78,7 +78,7 @@ class SemaOpenACC : public SemaBase { /// Collapse has an 'N' count that makes it apply to a number of loops 'below' /// it. struct CollapseCheckingInfo { - OpenACCCollapseClause *ActiveCollapse = nullptr; + const OpenACCCollapseClause *ActiveCollapse = nullptr; /// This is a value that maintains the current value of the 'N' on the /// current collapse, minus the depth that has already been traversed. When @@ -107,7 +107,7 @@ class SemaOpenACC : public SemaBase { /// own counting of elements. UnsignedOrNone CurTileCount = std::nullopt; - /// Records whether we've hit a 'CurTileCount' of '0' on the wya down, + /// Records whether we've hit a 'CurTileCount' of '0' on the way down, /// which allows us to diagnose if the number of arguments is too large for /// the current number of 'for' loops. bool TileDepthSatisfied = true; @@ -177,6 +177,21 @@ class SemaOpenACC : public SemaBase { void CheckLastRoutineDeclNameConflict(const NamedDecl *ND); + bool DiagnoseRequiredClauses(OpenACCDirectiveKind DK, SourceLocation DirLoc, + ArrayRef Clauses); + + bool DiagnoseAllowedClauses(OpenACCDirectiveKind DK, OpenACCClauseKind CK, + SourceLocation ClauseLoc); + +public: + // Needed from the visitor, so should be public. + bool DiagnoseAllowedOnceClauses(OpenACCDirectiveKind DK, OpenACCClauseKind CK, + SourceLocation ClauseLoc, + ArrayRef Clauses); + bool DiagnoseExclusiveClauses(OpenACCDirectiveKind DK, OpenACCClauseKind CK, + SourceLocation ClauseLoc, + ArrayRef Clauses); + public: ComputeConstructInfo &getActiveComputeConstructInfo() { return ActiveComputeConstructInfo; diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp index 2820d7b288658..01e183051934d 100644 --- a/clang/lib/AST/OpenACCClause.cpp +++ b/clang/lib/AST/OpenACCClause.cpp @@ -190,18 +190,15 @@ OpenACCCollapseClause::OpenACCCollapseClause(SourceLocation BeginLoc, SourceLocation EndLoc) : OpenACCClauseWithSingleIntExpr(OpenACCClauseKind::Collapse, BeginLoc, LParenLoc, LoopCount, EndLoc), - HasForce(HasForce) { - assert(LoopCount && "LoopCount required"); -} + HasForce(HasForce) {} OpenACCCollapseClause * OpenACCCollapseClause::Create(const ASTContext &C, SourceLocation BeginLoc, SourceLocation LParenLoc, bool HasForce, Expr *LoopCount, SourceLocation EndLoc) { - assert( - LoopCount && - (LoopCount->isInstantiationDependent() || isa(LoopCount)) && - "Loop count not constant expression"); + assert((!LoopCount || (LoopCount->isInstantiationDependent() || + isa(LoopCount))) && + "Loop count not constant expression"); void *Mem = C.Allocate(sizeof(OpenACCCollapseClause), alignof(OpenACCCollapseClause)); return new (Mem) diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp index 5f8417ac57b8b..51f7d17be4158 100644 --- a/clang/lib/AST/TextNodeDumper.cpp +++ b/clang/lib/AST/TextNodeDumper.cpp @@ -433,6 +433,7 @@ void TextNodeDumper::Visit(const OpenACCClause *C) { case OpenACCClauseKind::Vector: case OpenACCClauseKind::VectorLength: case OpenACCClauseKind::Invalid: + case OpenACCClauseKind::Shortloop: // The condition expression will be printed as a part of the 'children', // but print 'clause' here so it is clear what is happening from the dump. OS << " clause"; diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp index 337b3eca49764..3843e1ad51cf5 100644 --- a/clang/lib/Parse/ParseOpenACC.cpp +++ b/clang/lib/Parse/ParseOpenACC.cpp @@ -537,6 +537,8 @@ ClauseParensKind getClauseParensKind(OpenACCDirectiveKind DirKind, case OpenACCClauseKind::Tile: return ClauseParensKind::Required; + case OpenACCClauseKind::Shortloop: + llvm_unreachable("Shortloop shouldn't be generated in clang"); case OpenACCClauseKind::Auto: case OpenACCClauseKind::Finalize: case OpenACCClauseKind::IfPresent: diff --git a/clang/lib/Sema/CMakeLists.txt b/clang/lib/Sema/CMakeLists.txt index cc7921fc32254..d0b0508eb263a 100644 --- a/clang/lib/Sema/CMakeLists.txt +++ b/clang/lib/Sema/CMakeLists.txt @@ -75,6 +75,7 @@ add_clang_library(clangSema SemaOpenACC.cpp SemaOpenACCAtomic.cpp SemaOpenACCClause.cpp + SemaOpenACCClauseAppertainment.cpp SemaOpenCL.cpp SemaOpenMP.cpp SemaOverload.cpp diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp index c80f5f848f60b..74eee9b13ac6e 100644 --- a/clang/lib/Sema/SemaOpenACC.cpp +++ b/clang/lib/Sema/SemaOpenACC.cpp @@ -221,6 +221,48 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII( } } +namespace { +// Given two collapse clauses, and the uninstanted version of the new one, +// return the 'best' one for the purposes of setting the collapse checking +// values. +const OpenACCCollapseClause * +getBestCollapseCandidate(const OpenACCCollapseClause *Old, + const OpenACCCollapseClause *New, + const OpenACCCollapseClause *UnInstNew) { + // If the loop count is nullptr, it is because instantiation failed, so this + // can't be the best one. + if (!New->getLoopCount()) + return Old; + + // If the loop-count had an error, than 'new' isn't a candidate. + if (!New->getLoopCount()) + return Old; + + // Don't consider uninstantiated ones, since we can't really check these. + if (New->getLoopCount()->isInstantiationDependent()) + return Old; + + // If this is an instantiation, and the old version wasn't instantation + // dependent, than nothing has changed and we've already done a diagnostic + // based on this one, so don't consider it. + if (UnInstNew && !UnInstNew->getLoopCount()->isInstantiationDependent()) + return Old; + + // New is now a valid candidate, so if there isn't an old one at this point, + // New is the only valid one. + if (!Old) + return New; + + // If the 'New' expression has a larger value than 'Old', then it is the new + // best candidate. + if (cast(Old->getLoopCount())->getResultAsAPSInt() < + cast(New->getLoopCount())->getResultAsAPSInt()) + return New; + + return Old; +} +} // namespace + void SemaOpenACC::AssociatedStmtRAII::SetCollapseInfoBeforeAssociatedStmt( ArrayRef UnInstClauses, ArrayRef Clauses) { @@ -228,43 +270,57 @@ void SemaOpenACC::AssociatedStmtRAII::SetCollapseInfoBeforeAssociatedStmt( // Reset this checking for loops that aren't covered in a RAII object. SemaRef.LoopInfo.CurLevelHasLoopAlready = false; SemaRef.CollapseInfo.CollapseDepthSatisfied = true; + SemaRef.CollapseInfo.CurCollapseCount = 0; SemaRef.TileInfo.TileDepthSatisfied = true; // We make sure to take an optional list of uninstantiated clauses, so that // we can check to make sure we don't 'double diagnose' in the event that - // the value of 'N' was not dependent in a template. We also ensure during - // Sema that there is only 1 collapse on each construct, so we can count on - // the fact that if both find a 'collapse', that they are the same one. - auto *CollapseClauseItr = - llvm::find_if(Clauses, llvm::IsaPred); - auto *UnInstCollapseClauseItr = + // the value of 'N' was not dependent in a template. Since we cannot count on + // there only being a single collapse clause, we count on the order to make + // sure get the matching ones, and we count on TreeTransform not removing + // these, even if loop-count instantiation failed. We can check the + // non-dependent ones right away, and realize that subsequent instantiation + // can only make it more specific. + + auto *UnInstClauseItr = llvm::find_if(UnInstClauses, llvm::IsaPred); - - if (Clauses.end() == CollapseClauseItr) - return; - - OpenACCCollapseClause *CollapseClause = - cast(*CollapseClauseItr); - - SemaRef.CollapseInfo.ActiveCollapse = CollapseClause; - Expr *LoopCount = CollapseClause->getLoopCount(); - - // If the loop count is still instantiation dependent, setting the depth - // counter isn't necessary, so return here. - if (!LoopCount || LoopCount->isInstantiationDependent()) - return; - - // Suppress diagnostics if we've done a 'transform' where the previous version - // wasn't dependent, meaning we already diagnosed it. - if (UnInstCollapseClauseItr != UnInstClauses.end() && - !cast(*UnInstCollapseClauseItr) - ->getLoopCount() - ->isInstantiationDependent()) + auto *ClauseItr = + llvm::find_if(Clauses, llvm::IsaPred); + const OpenACCCollapseClause *FoundClause = nullptr; + + // Loop through the list of Collapse clauses and find the one that: + // 1- Has a non-dependent, non-null loop count (null means error, likely + // during instantiation). + // 2- If UnInstClauses isn't empty, its corresponding + // loop count was dependent. + // 3- Has the largest 'loop count' of all. + while (ClauseItr != Clauses.end()) { + const OpenACCCollapseClause *CurClause = + cast(*ClauseItr); + const OpenACCCollapseClause *UnInstCurClause = + UnInstClauseItr == UnInstClauses.end() + ? nullptr + : cast(*UnInstClauseItr); + + FoundClause = + getBestCollapseCandidate(FoundClause, CurClause, UnInstCurClause); + + UnInstClauseItr = + UnInstClauseItr == UnInstClauses.end() + ? UnInstClauseItr + : std::find_if(std::next(UnInstClauseItr), UnInstClauses.end(), + llvm::IsaPred); + ClauseItr = std::find_if(std::next(ClauseItr), Clauses.end(), + llvm::IsaPred); + } + + if (!FoundClause) return; + SemaRef.CollapseInfo.ActiveCollapse = FoundClause; SemaRef.CollapseInfo.CollapseDepthSatisfied = false; SemaRef.CollapseInfo.CurCollapseCount = - cast(LoopCount)->getResultAsAPSInt(); + cast(FoundClause->getLoopCount())->getResultAsAPSInt(); SemaRef.CollapseInfo.DirectiveKind = DirKind; } @@ -283,6 +339,17 @@ void SemaOpenACC::AssociatedStmtRAII::SetTileInfoBeforeAssociatedStmt( return; OpenACCTileClause *TileClause = cast(*TileClauseItr); + + // Multiple tile clauses are allowed, so ensure that we use the one with the + // largest 'tile count'. + while (Clauses.end() != + (TileClauseItr = std::find_if(std::next(TileClauseItr), Clauses.end(), + llvm::IsaPred))) { + OpenACCTileClause *NewClause = cast(*TileClauseItr); + if (NewClause->getSizeExprs().size() > TileClause->getSizeExprs().size()) + TileClause = NewClause; + } + SemaRef.TileInfo.ActiveTile = TileClause; SemaRef.TileInfo.TileDepthSatisfied = false; SemaRef.TileInfo.CurTileCount = @@ -1426,30 +1493,6 @@ void SemaOpenACC::ActOnForStmtEnd(SourceLocation ForLoc, StmtResult Body) { } namespace { -// Get a list of clause Kinds for diagnosing a list, joined by a commas and an -// 'or'. -std::string GetListOfClauses(llvm::ArrayRef Clauses) { - assert(!Clauses.empty() && "empty clause list not supported"); - - std::string Output; - llvm::raw_string_ostream OS{Output}; - - if (Clauses.size() == 1) { - OS << '\'' << Clauses[0] << '\''; - return Output; - } - - llvm::ArrayRef AllButLast{Clauses.begin(), - Clauses.end() - 1}; - - llvm::interleave( - AllButLast, [&](OpenACCClauseKind K) { OS << '\'' << K << '\''; }, - [&] { OS << ", "; }); - - OS << " or \'" << Clauses.back() << '\''; - return Output; -} - // Helper that should mirror ActOnRoutineName to get the FunctionDecl out for // magic-static checking. FunctionDecl *getFunctionFromRoutineName(Expr *RoutineName) { @@ -1682,86 +1725,8 @@ bool SemaOpenACC::ActOnStartStmtDirective( << OpenACCClauseKind::Tile; } - // OpenACC3.3 2.6.5: At least one copy, copyin, copyout, create, no_create, - // present, deviceptr, attach, or default clause must appear on a 'data' - // construct. - if (K == OpenACCDirectiveKind::Data && - llvm::find_if(Clauses, - llvm::IsaPred) == Clauses.end()) - return Diag(StartLoc, diag::err_acc_construct_one_clause_of) - << K - << GetListOfClauses( - {OpenACCClauseKind::Copy, OpenACCClauseKind::CopyIn, - OpenACCClauseKind::CopyOut, OpenACCClauseKind::Create, - OpenACCClauseKind::NoCreate, OpenACCClauseKind::Present, - OpenACCClauseKind::DevicePtr, OpenACCClauseKind::Attach, - OpenACCClauseKind::Default}); - - // OpenACC3.3 2.6.6: At least one copyin, create, or attach clause must appear - // on an enter data directive. - if (K == OpenACCDirectiveKind::EnterData && - llvm::find_if(Clauses, - llvm::IsaPred) == Clauses.end()) - return Diag(StartLoc, diag::err_acc_construct_one_clause_of) - << K - << GetListOfClauses({ - OpenACCClauseKind::CopyIn, - OpenACCClauseKind::Create, - OpenACCClauseKind::Attach, - }); - // OpenACC3.3 2.6.6: At least one copyout, delete, or detach clause must - // appear on an exit data directive. - if (K == OpenACCDirectiveKind::ExitData && - llvm::find_if(Clauses, - llvm::IsaPred) == Clauses.end()) - return Diag(StartLoc, diag::err_acc_construct_one_clause_of) - << K - << GetListOfClauses({ - OpenACCClauseKind::CopyOut, - OpenACCClauseKind::Delete, - OpenACCClauseKind::Detach, - }); - - // OpenACC3.3 2.8: At least 'one use_device' clause must appear. - if (K == OpenACCDirectiveKind::HostData && - llvm::find_if(Clauses, llvm::IsaPred) == - Clauses.end()) - return Diag(StartLoc, diag::err_acc_construct_one_clause_of) - << K << GetListOfClauses({OpenACCClauseKind::UseDevice}); - - // OpenACC3.3 2.14.3: At least one default_async, device_num, or device_type - // clause must appear. - if (K == OpenACCDirectiveKind::Set && - llvm::find_if( - Clauses, - llvm::IsaPred) == - Clauses.end()) - return Diag(StartLoc, diag::err_acc_construct_one_clause_of) - << K - << GetListOfClauses({OpenACCClauseKind::DefaultAsync, - OpenACCClauseKind::DeviceNum, - OpenACCClauseKind::DeviceType, - OpenACCClauseKind::If}); - - // OpenACC3.3 2.14.4: At least one self, host, or device clause must appear on - // an update directive. - if (K == OpenACCDirectiveKind::Update && - llvm::find_if(Clauses, llvm::IsaPred) == - Clauses.end()) - return Diag(StartLoc, diag::err_acc_construct_one_clause_of) - << K - << GetListOfClauses({OpenACCClauseKind::Self, - OpenACCClauseKind::Host, - OpenACCClauseKind::Device}); - + if (DiagnoseRequiredClauses(K, StartLoc, Clauses)) + return true; return diagnoseConstructAppertainment(*this, K, StartLoc, /*IsStmt=*/true); } @@ -1930,6 +1895,62 @@ StmtResult SemaOpenACC::ActOnAssociatedStmt( llvm_unreachable("Invalid associated statement application"); } +namespace { + +// Routine has some pretty complicated set of rules for how device_type +// interacts with 'gang', 'worker', 'vector', and 'seq'. Enforce part of it +// here. +bool CheckValidRoutineGangWorkerVectorSeqClauses( + SemaOpenACC &SemaRef, SourceLocation DirectiveLoc, + ArrayRef Clauses) { + auto RequiredPred = llvm::IsaPred; + // The clause handling has assured us that there is no duplicates. That is, + // if there is 1 before a device_type, there are none after a device_type. + // If not, there is at most 1 applying to each device_type. + + // What is left to legalize is that either: + // 1- there is 1 before the first device_type. + // 2- there is 1 AFTER each device_type. + auto *FirstDeviceType = + llvm::find_if(Clauses, llvm::IsaPred); + + // If there is 1 before the first device_type (or at all if no device_type), + // we are legal. + auto *ClauseItr = + std::find_if(Clauses.begin(), FirstDeviceType, RequiredPred); + + if (ClauseItr != FirstDeviceType) + return false; + + // If there IS no device_type, and no clause, diagnose. + if (FirstDeviceType == Clauses.end()) + return SemaRef.Diag(DirectiveLoc, diag::err_acc_construct_one_clause_of) + << OpenACCDirectiveKind::Routine + << "'gang', 'seq', 'vector', or 'worker'"; + + // Else, we have to check EACH device_type group. PrevDeviceType is the + // device-type before the current group. + auto *PrevDeviceType = FirstDeviceType; + + while (PrevDeviceType != Clauses.end()) { + auto *NextDeviceType = + std::find_if(std::next(PrevDeviceType), Clauses.end(), + llvm::IsaPred); + + ClauseItr = std::find_if(PrevDeviceType, NextDeviceType, RequiredPred); + + if (ClauseItr == NextDeviceType) + return SemaRef.Diag((*PrevDeviceType)->getBeginLoc(), + diag::err_acc_clause_routine_one_of_in_region); + + PrevDeviceType = NextDeviceType; + } + + return false; +} +} // namespace + bool SemaOpenACC::ActOnStartDeclDirective( OpenACCDirectiveKind K, SourceLocation StartLoc, ArrayRef Clauses) { @@ -1939,19 +1960,11 @@ bool SemaOpenACC::ActOnStartDeclDirective( SemaRef.DiscardCleanupsInEvaluationContext(); SemaRef.PopExpressionEvaluationContext(); + if (DiagnoseRequiredClauses(K, StartLoc, Clauses)) + return true; if (K == OpenACCDirectiveKind::Routine && - llvm::find_if(Clauses, - llvm::IsaPred) == - Clauses.end()) - return Diag(StartLoc, diag::err_acc_construct_one_clause_of) - << K - << GetListOfClauses({ - OpenACCClauseKind::Gang, - OpenACCClauseKind::Worker, - OpenACCClauseKind::Vector, - OpenACCClauseKind::Seq, - }); + CheckValidRoutineGangWorkerVectorSeqClauses(*this, StartLoc, Clauses)) + return true; return diagnoseConstructAppertainment(*this, K, StartLoc, /*IsStmt=*/false); } diff --git a/clang/lib/Sema/SemaOpenACCClause.cpp b/clang/lib/Sema/SemaOpenACCClause.cpp index 049baead031a1..3694a831b76de 100644 --- a/clang/lib/Sema/SemaOpenACCClause.cpp +++ b/clang/lib/Sema/SemaOpenACCClause.cpp @@ -20,502 +20,6 @@ using namespace clang; namespace { -bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind, - OpenACCClauseKind ClauseKind) { - switch (ClauseKind) { - // FIXME: For each clause as we implement them, we can add the - // 'legalization' list here. - case OpenACCClauseKind::Default: - switch (DirectiveKind) { - case OpenACCDirectiveKind::Parallel: - case OpenACCDirectiveKind::Serial: - case OpenACCDirectiveKind::Kernels: - case OpenACCDirectiveKind::ParallelLoop: - case OpenACCDirectiveKind::SerialLoop: - case OpenACCDirectiveKind::KernelsLoop: - case OpenACCDirectiveKind::Data: - return true; - default: - return false; - } - case OpenACCClauseKind::If: - switch (DirectiveKind) { - case OpenACCDirectiveKind::Parallel: - case OpenACCDirectiveKind::Serial: - case OpenACCDirectiveKind::Kernels: - case OpenACCDirectiveKind::Data: - case OpenACCDirectiveKind::EnterData: - case OpenACCDirectiveKind::ExitData: - case OpenACCDirectiveKind::HostData: - case OpenACCDirectiveKind::Init: - case OpenACCDirectiveKind::Shutdown: - case OpenACCDirectiveKind::Set: - case OpenACCDirectiveKind::Update: - case OpenACCDirectiveKind::Wait: - case OpenACCDirectiveKind::ParallelLoop: - case OpenACCDirectiveKind::SerialLoop: - case OpenACCDirectiveKind::KernelsLoop: - // OpenACC 3.4(prerelease) PR #511 adds 'if' to atomic. - case OpenACCDirectiveKind::Atomic: - return true; - default: - return false; - } - case OpenACCClauseKind::Self: - switch (DirectiveKind) { - case OpenACCDirectiveKind::Parallel: - case OpenACCDirectiveKind::Serial: - case OpenACCDirectiveKind::Kernels: - case OpenACCDirectiveKind::Update: - case OpenACCDirectiveKind::ParallelLoop: - case OpenACCDirectiveKind::SerialLoop: - case OpenACCDirectiveKind::KernelsLoop: - return true; - default: - return false; - } - case OpenACCClauseKind::NumGangs: - case OpenACCClauseKind::NumWorkers: - case OpenACCClauseKind::VectorLength: - switch (DirectiveKind) { - case OpenACCDirectiveKind::Parallel: - case OpenACCDirectiveKind::Kernels: - case OpenACCDirectiveKind::ParallelLoop: - case OpenACCDirectiveKind::KernelsLoop: - return true; - default: - return false; - } - case OpenACCClauseKind::FirstPrivate: - switch (DirectiveKind) { - case OpenACCDirectiveKind::Parallel: - case OpenACCDirectiveKind::Serial: - case OpenACCDirectiveKind::ParallelLoop: - case OpenACCDirectiveKind::SerialLoop: - return true; - default: - return false; - } - case OpenACCClauseKind::Private: - switch (DirectiveKind) { - case OpenACCDirectiveKind::Parallel: - case OpenACCDirectiveKind::Serial: - case OpenACCDirectiveKind::Loop: - case OpenACCDirectiveKind::ParallelLoop: - case OpenACCDirectiveKind::SerialLoop: - case OpenACCDirectiveKind::KernelsLoop: - return true; - default: - return false; - } - case OpenACCClauseKind::NoCreate: - switch (DirectiveKind) { - case OpenACCDirectiveKind::Parallel: - case OpenACCDirectiveKind::Serial: - case OpenACCDirectiveKind::Kernels: - case OpenACCDirectiveKind::Data: - case OpenACCDirectiveKind::ParallelLoop: - case OpenACCDirectiveKind::SerialLoop: - case OpenACCDirectiveKind::KernelsLoop: - return true; - default: - return false; - } - case OpenACCClauseKind::Present: - switch (DirectiveKind) { - case OpenACCDirectiveKind::Parallel: - case OpenACCDirectiveKind::Serial: - case OpenACCDirectiveKind::Kernels: - case OpenACCDirectiveKind::Data: - case OpenACCDirectiveKind::Declare: - case OpenACCDirectiveKind::ParallelLoop: - case OpenACCDirectiveKind::SerialLoop: - case OpenACCDirectiveKind::KernelsLoop: - return true; - default: - return false; - } - - case OpenACCClauseKind::Copy: - case OpenACCClauseKind::PCopy: - case OpenACCClauseKind::PresentOrCopy: - switch (DirectiveKind) { - case OpenACCDirectiveKind::Parallel: - case OpenACCDirectiveKind::Serial: - case OpenACCDirectiveKind::Kernels: - case OpenACCDirectiveKind::Data: - case OpenACCDirectiveKind::Declare: - case OpenACCDirectiveKind::ParallelLoop: - case OpenACCDirectiveKind::SerialLoop: - case OpenACCDirectiveKind::KernelsLoop: - return true; - default: - return false; - } - case OpenACCClauseKind::CopyIn: - case OpenACCClauseKind::PCopyIn: - case OpenACCClauseKind::PresentOrCopyIn: - switch (DirectiveKind) { - case OpenACCDirectiveKind::Parallel: - case OpenACCDirectiveKind::Serial: - case OpenACCDirectiveKind::Kernels: - case OpenACCDirectiveKind::Data: - case OpenACCDirectiveKind::EnterData: - case OpenACCDirectiveKind::Declare: - case OpenACCDirectiveKind::ParallelLoop: - case OpenACCDirectiveKind::SerialLoop: - case OpenACCDirectiveKind::KernelsLoop: - return true; - default: - return false; - } - case OpenACCClauseKind::CopyOut: - case OpenACCClauseKind::PCopyOut: - case OpenACCClauseKind::PresentOrCopyOut: - switch (DirectiveKind) { - case OpenACCDirectiveKind::Parallel: - case OpenACCDirectiveKind::Serial: - case OpenACCDirectiveKind::Kernels: - case OpenACCDirectiveKind::Data: - case OpenACCDirectiveKind::ExitData: - case OpenACCDirectiveKind::Declare: - case OpenACCDirectiveKind::ParallelLoop: - case OpenACCDirectiveKind::SerialLoop: - case OpenACCDirectiveKind::KernelsLoop: - return true; - default: - return false; - } - case OpenACCClauseKind::Create: - case OpenACCClauseKind::PCreate: - case OpenACCClauseKind::PresentOrCreate: - switch (DirectiveKind) { - case OpenACCDirectiveKind::Parallel: - case OpenACCDirectiveKind::Serial: - case OpenACCDirectiveKind::Kernels: - case OpenACCDirectiveKind::Data: - case OpenACCDirectiveKind::EnterData: - case OpenACCDirectiveKind::Declare: - case OpenACCDirectiveKind::ParallelLoop: - case OpenACCDirectiveKind::SerialLoop: - case OpenACCDirectiveKind::KernelsLoop: - return true; - default: - return false; - } - - case OpenACCClauseKind::Attach: - switch (DirectiveKind) { - case OpenACCDirectiveKind::Parallel: - case OpenACCDirectiveKind::Serial: - case OpenACCDirectiveKind::Kernels: - case OpenACCDirectiveKind::Data: - case OpenACCDirectiveKind::EnterData: - case OpenACCDirectiveKind::ParallelLoop: - case OpenACCDirectiveKind::SerialLoop: - case OpenACCDirectiveKind::KernelsLoop: - return true; - default: - return false; - } - case OpenACCClauseKind::DevicePtr: - switch (DirectiveKind) { - case OpenACCDirectiveKind::Parallel: - case OpenACCDirectiveKind::Serial: - case OpenACCDirectiveKind::Kernels: - case OpenACCDirectiveKind::Data: - case OpenACCDirectiveKind::Declare: - case OpenACCDirectiveKind::ParallelLoop: - case OpenACCDirectiveKind::SerialLoop: - case OpenACCDirectiveKind::KernelsLoop: - return true; - default: - return false; - } - case OpenACCClauseKind::Async: - switch (DirectiveKind) { - case OpenACCDirectiveKind::Parallel: - case OpenACCDirectiveKind::Serial: - case OpenACCDirectiveKind::Kernels: - case OpenACCDirectiveKind::Data: - case OpenACCDirectiveKind::EnterData: - case OpenACCDirectiveKind::ExitData: - case OpenACCDirectiveKind::Set: - case OpenACCDirectiveKind::Update: - case OpenACCDirectiveKind::Wait: - case OpenACCDirectiveKind::ParallelLoop: - case OpenACCDirectiveKind::SerialLoop: - case OpenACCDirectiveKind::KernelsLoop: - return true; - default: - return false; - } - case OpenACCClauseKind::Wait: - switch (DirectiveKind) { - case OpenACCDirectiveKind::Parallel: - case OpenACCDirectiveKind::Serial: - case OpenACCDirectiveKind::Kernels: - case OpenACCDirectiveKind::Data: - case OpenACCDirectiveKind::EnterData: - case OpenACCDirectiveKind::ExitData: - case OpenACCDirectiveKind::Update: - case OpenACCDirectiveKind::ParallelLoop: - case OpenACCDirectiveKind::SerialLoop: - case OpenACCDirectiveKind::KernelsLoop: - return true; - default: - return false; - } - - case OpenACCClauseKind::Seq: - switch (DirectiveKind) { - case OpenACCDirectiveKind::Loop: - case OpenACCDirectiveKind::Routine: - case OpenACCDirectiveKind::ParallelLoop: - case OpenACCDirectiveKind::SerialLoop: - case OpenACCDirectiveKind::KernelsLoop: - return true; - default: - return false; - } - - case OpenACCClauseKind::Independent: - case OpenACCClauseKind::Auto: - switch (DirectiveKind) { - case OpenACCDirectiveKind::Loop: - case OpenACCDirectiveKind::ParallelLoop: - case OpenACCDirectiveKind::SerialLoop: - case OpenACCDirectiveKind::KernelsLoop: - return true; - default: - return false; - } - - case OpenACCClauseKind::Reduction: - switch (DirectiveKind) { - case OpenACCDirectiveKind::Parallel: - case OpenACCDirectiveKind::Serial: - case OpenACCDirectiveKind::Loop: - case OpenACCDirectiveKind::ParallelLoop: - case OpenACCDirectiveKind::SerialLoop: - case OpenACCDirectiveKind::KernelsLoop: - return true; - default: - return false; - } - - case OpenACCClauseKind::DeviceType: - case OpenACCClauseKind::DType: - switch (DirectiveKind) { - case OpenACCDirectiveKind::Parallel: - case OpenACCDirectiveKind::Serial: - case OpenACCDirectiveKind::Kernels: - case OpenACCDirectiveKind::Data: - case OpenACCDirectiveKind::Init: - case OpenACCDirectiveKind::Shutdown: - case OpenACCDirectiveKind::Set: - case OpenACCDirectiveKind::Update: - case OpenACCDirectiveKind::Loop: - case OpenACCDirectiveKind::Routine: - case OpenACCDirectiveKind::ParallelLoop: - case OpenACCDirectiveKind::SerialLoop: - case OpenACCDirectiveKind::KernelsLoop: - return true; - default: - return false; - } - - case OpenACCClauseKind::Collapse: { - switch (DirectiveKind) { - case OpenACCDirectiveKind::Loop: - case OpenACCDirectiveKind::ParallelLoop: - case OpenACCDirectiveKind::SerialLoop: - case OpenACCDirectiveKind::KernelsLoop: - return true; - default: - return false; - } - } - case OpenACCClauseKind::Tile: { - switch (DirectiveKind) { - case OpenACCDirectiveKind::Loop: - case OpenACCDirectiveKind::ParallelLoop: - case OpenACCDirectiveKind::SerialLoop: - case OpenACCDirectiveKind::KernelsLoop: - return true; - default: - return false; - } - } - - case OpenACCClauseKind::Gang: { - switch (DirectiveKind) { - case OpenACCDirectiveKind::Loop: - case OpenACCDirectiveKind::ParallelLoop: - case OpenACCDirectiveKind::SerialLoop: - case OpenACCDirectiveKind::KernelsLoop: - case OpenACCDirectiveKind::Routine: - return true; - default: - return false; - } - case OpenACCClauseKind::Worker: { - switch (DirectiveKind) { - case OpenACCDirectiveKind::Loop: - case OpenACCDirectiveKind::ParallelLoop: - case OpenACCDirectiveKind::SerialLoop: - case OpenACCDirectiveKind::KernelsLoop: - case OpenACCDirectiveKind::Routine: - return true; - default: - return false; - } - } - case OpenACCClauseKind::Vector: { - switch (DirectiveKind) { - case OpenACCDirectiveKind::Loop: - case OpenACCDirectiveKind::ParallelLoop: - case OpenACCDirectiveKind::SerialLoop: - case OpenACCDirectiveKind::KernelsLoop: - case OpenACCDirectiveKind::Routine: - return true; - default: - return false; - } - } - case OpenACCClauseKind::Finalize: { - switch (DirectiveKind) { - case OpenACCDirectiveKind::ExitData: - return true; - default: - return false; - } - } - case OpenACCClauseKind::IfPresent: { - switch (DirectiveKind) { - case OpenACCDirectiveKind::HostData: - case OpenACCDirectiveKind::Update: - return true; - default: - return false; - } - } - case OpenACCClauseKind::Delete: { - switch (DirectiveKind) { - case OpenACCDirectiveKind::ExitData: - return true; - default: - return false; - } - } - - case OpenACCClauseKind::Detach: { - switch (DirectiveKind) { - case OpenACCDirectiveKind::ExitData: - return true; - default: - return false; - } - } - - case OpenACCClauseKind::DeviceNum: { - switch (DirectiveKind) { - case OpenACCDirectiveKind::Init: - case OpenACCDirectiveKind::Shutdown: - case OpenACCDirectiveKind::Set: - return true; - default: - return false; - } - } - case OpenACCClauseKind::Link: { - switch (DirectiveKind) { - case OpenACCDirectiveKind::Declare: - return true; - default: - return false; - } - } - case OpenACCClauseKind::DeviceResident: { - switch (DirectiveKind) { - case OpenACCDirectiveKind::Declare: - return true; - default: - return false; - } - } - - case OpenACCClauseKind::UseDevice: { - switch (DirectiveKind) { - case OpenACCDirectiveKind::HostData: - return true; - default: - return false; - } - } - case OpenACCClauseKind::DefaultAsync: { - switch (DirectiveKind) { - case OpenACCDirectiveKind::Set: - return true; - default: - return false; - } - } - case OpenACCClauseKind::Device: { - switch (DirectiveKind) { - case OpenACCDirectiveKind::Update: - return true; - default: - return false; - } - } - case OpenACCClauseKind::Host: { - switch (DirectiveKind) { - case OpenACCDirectiveKind::Update: - return true; - default: - return false; - } - } - case OpenACCClauseKind::NoHost: { - switch (DirectiveKind) { - case OpenACCDirectiveKind::Routine: - return true; - default: - return false; - } - } - case OpenACCClauseKind::Bind: { - switch (DirectiveKind) { - case OpenACCDirectiveKind::Routine: - return true; - default: - return false; - } - } - } - - default: - // Do nothing so we can go to the 'unimplemented' diagnostic instead. - return true; - } - llvm_unreachable("Invalid clause kind"); -} - -bool checkAlreadyHasClauseOfKind( - SemaOpenACC &S, ArrayRef ExistingClauses, - SemaOpenACC::OpenACCParsedClause &Clause) { - const auto *Itr = llvm::find_if(ExistingClauses, [&](const OpenACCClause *C) { - return C->getClauseKind() == Clause.getClauseKind(); - }); - if (Itr != ExistingClauses.end()) { - S.Diag(Clause.getBeginLoc(), diag::err_acc_duplicate_clause_disallowed) - << Clause.getDirectiveKind() << Clause.getClauseKind(); - S.Diag((*Itr)->getBeginLoc(), diag::note_acc_previous_clause_here); - return true; - } - return false; -} bool checkValidAfterDeviceType( SemaOpenACC &S, const OpenACCDeviceTypeClause &DeviceTypeClause, const SemaOpenACC::OpenACCParsedClause &NewClause) { @@ -627,14 +131,6 @@ bool checkValidAfterDeviceType( return true; } -// A temporary function that helps implement the 'not implemented' check at the -// top of each clause checking function. This should only be used in conjunction -// with the one being currently implemented/only updated after the entire -// construct has been implemented. -bool isDirectiveKindImplemented(OpenACCDirectiveKind DK) { - return DK != OpenACCDirectiveKind::Routine; -} - // GCC looks through linkage specs, but not the other transparent declaration // contexts for 'declare' restrictions, so this helper function helps get us // through that. @@ -649,28 +145,20 @@ class SemaOpenACCClauseVisitor { SemaOpenACC &SemaRef; ASTContext &Ctx; ArrayRef ExistingClauses; - bool NotImplemented = false; - - OpenACCClause *isNotImplemented() { - NotImplemented = true; - return nullptr; - } // OpenACC 3.3 2.9: - // A 'gang', 'worker', or 'vector' clause may not appear if a 'seq' clause - // appears. - // -also- - // OpenACC3.3 2.15: (routine) - // Exactly one of the 'gang', 'worker', 'vector' or 'seq' clauses must appear. + // A 'gang', 'worker', or 'vector' clause may not appear if a 'seq' clause + // appears. bool DiagGangWorkerVectorSeqConflict(SemaOpenACC::OpenACCParsedClause &Clause) { + if (Clause.getDirectiveKind() != OpenACCDirectiveKind::Loop && + !isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())) + return false; + assert(Clause.getClauseKind() == OpenACCClauseKind::Gang || + Clause.getClauseKind() == OpenACCClauseKind::Worker || + Clause.getClauseKind() == OpenACCClauseKind::Vector); const auto *Itr = - Clause.getDirectiveKind() == OpenACCDirectiveKind::Routine - ? llvm::find_if( - ExistingClauses, - llvm::IsaPred) - : llvm::find_if(ExistingClauses, llvm::IsaPred); + llvm::find_if(ExistingClauses, llvm::IsaPred); if (Itr != ExistingClauses.end()) { SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_cannot_combine) @@ -736,13 +224,165 @@ class SemaOpenACCClauseVisitor { llvm_unreachable("didn't return from switch above?"); } + // Helper for the 'routine' checks during 'new' clause addition. Precondition + // is that we already know the new clause is one of the prohbiited ones. + template + bool + CheckValidRoutineNewClauseHelper(Pred HasPredicate, + SemaOpenACC::OpenACCParsedClause &Clause) { + if (Clause.getDirectiveKind() != OpenACCDirectiveKind::Routine) + return false; + + auto *FirstDeviceType = + llvm::find_if(ExistingClauses, llvm::IsaPred); + + if (FirstDeviceType == ExistingClauses.end()) { + // If there isn't a device type yet, ANY duplicate is wrong. + + auto *ExistingProhibitedClause = + llvm::find_if(ExistingClauses, HasPredicate); + + if (ExistingProhibitedClause == ExistingClauses.end()) + return false; + + SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_cannot_combine) + << Clause.getClauseKind() + << (*ExistingProhibitedClause)->getClauseKind() + << Clause.getDirectiveKind(); + SemaRef.Diag((*ExistingProhibitedClause)->getBeginLoc(), + diag::note_acc_previous_clause_here); + return true; + } + + // At this point we know that this is 'after' a device type. So this is an + // error if: 1- there is one BEFORE the 'device_type' 2- there is one + // between this and the previous 'device_type'. + + auto *BeforeDeviceType = + std::find_if(ExistingClauses.begin(), FirstDeviceType, HasPredicate); + // If there is one before the device_type (and we know we are after a + // device_type), than this is ill-formed. + if (BeforeDeviceType != FirstDeviceType) { + SemaRef.Diag( + Clause.getBeginLoc(), + diag::err_acc_clause_routine_cannot_combine_before_device_type) + << Clause.getClauseKind() << (*BeforeDeviceType)->getClauseKind(); + SemaRef.Diag((*BeforeDeviceType)->getBeginLoc(), + diag::note_acc_previous_clause_here); + SemaRef.Diag((*FirstDeviceType)->getBeginLoc(), + diag::note_acc_previous_clause_here); + return true; + } + + auto LastDeviceTypeItr = + std::find_if(ExistingClauses.rbegin(), ExistingClauses.rend(), + llvm::IsaPred); + + // We already know there is one in the list, so it is nonsensical to not + // have one. + assert(LastDeviceTypeItr != ExistingClauses.rend()); + + // Get the device-type from-the-front (not reverse) iterator from the + // reverse iterator. + auto *LastDeviceType = LastDeviceTypeItr.base() - 1; + + auto *ExistingProhibitedSinceLastDevice = + std::find_if(LastDeviceType, ExistingClauses.end(), HasPredicate); + + // No prohibited ones since the last device-type. + if (ExistingProhibitedSinceLastDevice == ExistingClauses.end()) + return false; + + SemaRef.Diag(Clause.getBeginLoc(), + diag::err_acc_clause_routine_cannot_combine_same_device_type) + << Clause.getClauseKind() + << (*ExistingProhibitedSinceLastDevice)->getClauseKind(); + SemaRef.Diag((*ExistingProhibitedSinceLastDevice)->getBeginLoc(), + diag::note_acc_previous_clause_here); + SemaRef.Diag((*LastDeviceType)->getBeginLoc(), + diag::note_acc_previous_clause_here); + return true; + } + + // Routine has a pretty complicated set of rules for how device_type and the + // gang, worker, vector, and seq clauses work. So diagnose some of it here. + bool CheckValidRoutineGangWorkerVectorSeqNewClause( + SemaOpenACC::OpenACCParsedClause &Clause) { + + if (Clause.getClauseKind() != OpenACCClauseKind::Gang && + Clause.getClauseKind() != OpenACCClauseKind::Vector && + Clause.getClauseKind() != OpenACCClauseKind::Worker && + Clause.getClauseKind() != OpenACCClauseKind::Seq) + return false; + auto ProhibitedPred = llvm::IsaPred; + + return CheckValidRoutineNewClauseHelper(ProhibitedPred, Clause); + } + + // Bind should have similar rules on a routine as gang/worker/vector/seq, + // except there is no 'must have 1' rule, so we can get all the checking done + // here. + bool + CheckValidRoutineBindNewClause(SemaOpenACC::OpenACCParsedClause &Clause) { + + if (Clause.getClauseKind() != OpenACCClauseKind::Bind) + return false; + + auto HasBindPred = llvm::IsaPred; + return CheckValidRoutineNewClauseHelper(HasBindPred, Clause); + } + + // For 'tile' and 'collapse', only allow 1 per 'device_type'. + // Also applies to num_worker, num_gangs, and vector_length. + template + bool DisallowSinceLastDeviceType(SemaOpenACC::OpenACCParsedClause &Clause) { + auto LastDeviceTypeItr = + std::find_if(ExistingClauses.rbegin(), ExistingClauses.rend(), + llvm::IsaPred); + + auto Last = std::find_if(ExistingClauses.rbegin(), LastDeviceTypeItr, + llvm::IsaPred); + + if (Last == LastDeviceTypeItr) + return false; + + SemaRef.Diag(Clause.getBeginLoc(), + diag::err_acc_clause_since_last_device_type) + << Clause.getClauseKind() << Clause.getDirectiveKind() + << (LastDeviceTypeItr != ExistingClauses.rend()); + SemaRef.Diag((*Last)->getBeginLoc(), diag::note_acc_previous_clause_here); + + if (LastDeviceTypeItr != ExistingClauses.rend()) + SemaRef.Diag((*LastDeviceTypeItr)->getBeginLoc(), + diag::note_acc_previous_clause_here); + + return true; + } + public: SemaOpenACCClauseVisitor(SemaOpenACC &S, ArrayRef ExistingClauses) : SemaRef(S), Ctx(S.getASTContext()), ExistingClauses(ExistingClauses) {} OpenACCClause *Visit(SemaOpenACC::OpenACCParsedClause &Clause) { + + if (SemaRef.DiagnoseAllowedOnceClauses( + Clause.getDirectiveKind(), Clause.getClauseKind(), + Clause.getBeginLoc(), ExistingClauses) || + SemaRef.DiagnoseExclusiveClauses(Clause.getDirectiveKind(), + Clause.getClauseKind(), + Clause.getBeginLoc(), ExistingClauses)) + return nullptr; + if (CheckValidRoutineGangWorkerVectorSeqNewClause(Clause) || + CheckValidRoutineBindNewClause(Clause)) + return nullptr; + switch (Clause.getClauseKind()) { + case OpenACCClauseKind::Shortloop: + llvm_unreachable("Shortloop shouldn't be generated in clang"); + case OpenACCClauseKind::Invalid: + return nullptr; #define VISIT_CLAUSE(CLAUSE_NAME) \ case OpenACCClauseKind::CLAUSE_NAME: \ return Visit##CLAUSE_NAME##Clause(Clause); @@ -753,8 +393,6 @@ class SemaOpenACCClauseVisitor { << Clause.getClauseKind() << OpenACCClauseKind::CLAUSE_NAME; \ return Visit##CLAUSE_NAME##Clause(Clause); #include "clang/Basic/OpenACCClauses.def" - default: - return isNotImplemented(); } llvm_unreachable("Invalid clause kind"); } @@ -771,13 +409,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitDefaultClause( if (Clause.getDefaultClauseKind() == OpenACCDefaultClauseKind::Invalid) return nullptr; - // OpenACC 3.3, Section 2.5.4: - // At most one 'default' clause may appear, and it must have a value of - // either 'none' or 'present'. - // Second half of the sentence is diagnosed during parsing. - if (checkAlreadyHasClauseOfKind(SemaRef, ExistingClauses, Clause)) - return nullptr; - return OpenACCDefaultClause::Create( Ctx, Clause.getDefaultClauseKind(), Clause.getBeginLoc(), Clause.getLParenLoc(), Clause.getEndLoc()); @@ -786,12 +417,7 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitDefaultClause( OpenACCClause *SemaOpenACCClauseVisitor::VisitTileClause( SemaOpenACC::OpenACCParsedClause &Clause) { - // Duplicates here are not really sensible. We could possible permit - // multiples if they all had the same value, but there isn't really a good - // reason to do so. Also, this simplifies the suppression of duplicates, in - // that we know if we 'find' one after instantiation, that it is the same - // clause, which simplifies instantiation/checking/etc. - if (checkAlreadyHasClauseOfKind(SemaRef, ExistingClauses, Clause)) + if (DisallowSinceLastDeviceType(Clause)) return nullptr; llvm::SmallVector NewSizeExprs; @@ -813,16 +439,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitTileClause( OpenACCClause *SemaOpenACCClauseVisitor::VisitIfClause( SemaOpenACC::OpenACCParsedClause &Clause) { - // There is no prose in the standard that says duplicates aren't allowed, - // but this diagnostic is present in other compilers, as well as makes - // sense. Prose DOES exist for 'data' and 'host_data', 'set', 'enter data' and - // 'exit data' both don't, but other implmementations do this. OpenACC issue - // 519 filed for the latter two. Prose also exists for 'update'. - // GCC allows this on init/shutdown, presumably for good reason, so we do too. - if (Clause.getDirectiveKind() != OpenACCDirectiveKind::Init && - Clause.getDirectiveKind() != OpenACCDirectiveKind::Shutdown && - checkAlreadyHasClauseOfKind(SemaRef, ExistingClauses, Clause)) - return nullptr; // The parser has ensured that we have a proper condition expr, so there // isn't really much to do here. @@ -845,11 +461,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitIfClause( OpenACCClause *SemaOpenACCClauseVisitor::VisitSelfClause( SemaOpenACC::OpenACCParsedClause &Clause) { - // There is no prose in the standard that says duplicates aren't allowed, - // but this diagnostic is present in other compilers, as well as makes - // sense. - if (checkAlreadyHasClauseOfKind(SemaRef, ExistingClauses, Clause)) - return nullptr; // If the 'if' clause is true, it makes the 'self' clause have no effect, // diagnose that here. This only applies on compute/combined constructs. @@ -871,10 +482,8 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitSelfClause( OpenACCClause *SemaOpenACCClauseVisitor::VisitNumGangsClause( SemaOpenACC::OpenACCParsedClause &Clause) { - // There is no prose in the standard that says duplicates aren't allowed, - // but this diagnostic is present in other compilers, as well as makes - // sense. - if (checkAlreadyHasClauseOfKind(SemaRef, ExistingClauses, Clause)) + + if (DisallowSinceLastDeviceType(Clause)) return nullptr; // num_gangs requires at least 1 int expr in all forms. Diagnose here, but @@ -966,10 +575,8 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitNumGangsClause( OpenACCClause *SemaOpenACCClauseVisitor::VisitNumWorkersClause( SemaOpenACC::OpenACCParsedClause &Clause) { - // There is no prose in the standard that says duplicates aren't allowed, - // but this diagnostic is present in other compilers, as well as makes - // sense. - if (checkAlreadyHasClauseOfKind(SemaRef, ExistingClauses, Clause)) + + if (DisallowSinceLastDeviceType(Clause)) return nullptr; // OpenACC 3.3 Section 2.9.2: @@ -1000,10 +607,8 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitNumWorkersClause( OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorLengthClause( SemaOpenACC::OpenACCParsedClause &Clause) { - // There is no prose in the standard that says duplicates aren't allowed, - // but this diagnostic is present in other compilers, as well as makes - // sense. - if (checkAlreadyHasClauseOfKind(SemaRef, ExistingClauses, Clause)) + + if (DisallowSinceLastDeviceType(Clause)) return nullptr; // OpenACC 3.3 Section 2.9.4: @@ -1034,12 +639,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorLengthClause( OpenACCClause *SemaOpenACCClauseVisitor::VisitAsyncClause( SemaOpenACC::OpenACCParsedClause &Clause) { - // There is no prose in the standard that says duplicates aren't allowed, - // but this diagnostic is present in other compilers, as well as makes - // sense. - if (checkAlreadyHasClauseOfKind(SemaRef, ExistingClauses, Clause)) - return nullptr; - assert(Clause.getNumIntExprs() < 2 && "Invalid number of expressions for Async"); return OpenACCAsyncClause::Create( @@ -1050,17 +649,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitAsyncClause( OpenACCClause *SemaOpenACCClauseVisitor::VisitDeviceNumClause( SemaOpenACC::OpenACCParsedClause &Clause) { - // Restrictions only properly implemented on certain constructs, so skip/treat - // as unimplemented in those cases. - if (!isDirectiveKindImplemented(Clause.getDirectiveKind())) - return isNotImplemented(); - - // OpenACC 3.3 2.14.3: Two instances of the same clause may not appear on the - // same directive. - if (Clause.getDirectiveKind() == OpenACCDirectiveKind::Set && - checkAlreadyHasClauseOfKind(SemaRef, ExistingClauses, Clause)) - return nullptr; - assert(Clause.getNumIntExprs() == 1 && "Invalid number of expressions for device_num"); return OpenACCDeviceNumClause::Create( @@ -1070,11 +658,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitDeviceNumClause( OpenACCClause *SemaOpenACCClauseVisitor::VisitDefaultAsyncClause( SemaOpenACC::OpenACCParsedClause &Clause) { - // OpenACC 3.3 2.14.3: Two instances of the same clause may not appear on the - // same directive. - if (checkAlreadyHasClauseOfKind(SemaRef, ExistingClauses, Clause)) - return nullptr; - assert(Clause.getNumIntExprs() == 1 && "Invalid number of expressions for default_async"); return OpenACCDefaultAsyncClause::Create( @@ -1332,11 +915,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitWaitClause( OpenACCClause *SemaOpenACCClauseVisitor::VisitDeviceTypeClause( SemaOpenACC::OpenACCParsedClause &Clause) { - // OpenACC 3.3 2.14.3: Two instances of the same clause may not appear on the - // same directive. - if (Clause.getDirectiveKind() == OpenACCDirectiveKind::Set && - checkAlreadyHasClauseOfKind(SemaRef, ExistingClauses, Clause)) - return nullptr; // Based on discussions, having more than 1 'architecture' on a 'set' is // nonsensical, so we're going to fix the standard to reflect this. Implement @@ -1397,18 +975,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitDeviceTypeClause( OpenACCClause *SemaOpenACCClauseVisitor::VisitAutoClause( SemaOpenACC::OpenACCParsedClause &Clause) { - // OpenACC 3.3 2.9: - // Only one of the seq, independent, and auto clauses may appear. - const auto *Itr = - llvm::find_if(ExistingClauses, - llvm::IsaPred); - if (Itr != ExistingClauses.end()) { - SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_loop_spec_conflict) - << Clause.getClauseKind() << Clause.getDirectiveKind(); - SemaRef.Diag((*Itr)->getBeginLoc(), diag::note_acc_previous_clause_here); - return nullptr; - } return OpenACCAutoClause::Create(Ctx, Clause.getBeginLoc(), Clause.getEndLoc()); @@ -1422,17 +988,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitNoHostClause( OpenACCClause *SemaOpenACCClauseVisitor::VisitIndependentClause( SemaOpenACC::OpenACCParsedClause &Clause) { - // OpenACC 3.3 2.9: - // Only one of the seq, independent, and auto clauses may appear. - const auto *Itr = llvm::find_if( - ExistingClauses, llvm::IsaPred); - if (Itr != ExistingClauses.end()) { - SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_loop_spec_conflict) - << Clause.getClauseKind() << Clause.getDirectiveKind(); - SemaRef.Diag((*Itr)->getBeginLoc(), diag::note_acc_previous_clause_here); - return nullptr; - } return OpenACCIndependentClause::Create(Ctx, Clause.getBeginLoc(), Clause.getEndLoc()); @@ -1805,6 +1360,7 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitWorkerClause( OpenACCClause *SemaOpenACCClauseVisitor::VisitGangClause( SemaOpenACC::OpenACCParsedClause &Clause) { + if (DiagGangWorkerVectorSeqConflict(Clause)) return nullptr; @@ -1944,45 +1500,23 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitIfPresentClause( OpenACCClause *SemaOpenACCClauseVisitor::VisitSeqClause( SemaOpenACC::OpenACCParsedClause &Clause) { - - if (Clause.getDirectiveKind() != OpenACCDirectiveKind::Routine) { - // OpenACC 3.3 2.9: - // Only one of the seq, independent, and auto clauses may appear. - const auto *Itr = - llvm::find_if(ExistingClauses, - llvm::IsaPred); + // OpenACC 3.3 2.9: + // A 'gang', 'worker', or 'vector' clause may not appear if a 'seq' clause + // appears. + if (Clause.getDirectiveKind() == OpenACCDirectiveKind::Loop || + isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())) { + const auto *Itr = llvm::find_if( + ExistingClauses, llvm::IsaPred); if (Itr != ExistingClauses.end()) { - SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_loop_spec_conflict) - << Clause.getClauseKind() << Clause.getDirectiveKind(); + SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_cannot_combine) + << Clause.getClauseKind() << (*Itr)->getClauseKind() + << Clause.getDirectiveKind(); SemaRef.Diag((*Itr)->getBeginLoc(), diag::note_acc_previous_clause_here); return nullptr; } } - // OpenACC 3.3 2.9: - // A 'gang', 'worker', or 'vector' clause may not appear if a 'seq' clause - // appears. - // -also- - // OpenACC3.3 2.15: (routine) - // Exactly one of the 'gang', 'worker', 'vector' or 'seq' clauses must appear. - const auto *Itr = - Clause.getDirectiveKind() == OpenACCDirectiveKind::Routine - ? llvm::find_if(ExistingClauses, - llvm::IsaPred) - : llvm::find_if(ExistingClauses, - llvm::IsaPred); - - if (Itr != ExistingClauses.end()) { - SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_cannot_combine) - << Clause.getClauseKind() << (*Itr)->getClauseKind() - << Clause.getDirectiveKind(); - SemaRef.Diag((*Itr)->getBeginLoc(), diag::note_acc_previous_clause_here); - return nullptr; - } - return OpenACCSeqClause::Create(Ctx, Clause.getBeginLoc(), Clause.getEndLoc()); } @@ -2083,12 +1617,8 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitReductionClause( OpenACCClause *SemaOpenACCClauseVisitor::VisitCollapseClause( SemaOpenACC::OpenACCParsedClause &Clause) { - // Duplicates here are not really sensible. We could possible permit - // multiples if they all had the same value, but there isn't really a good - // reason to do so. Also, this simplifies the suppression of duplicates, in - // that we know if we 'find' one after instantiation, that it is the same - // clause, which simplifies instantiation/checking/etc. - if (checkAlreadyHasClauseOfKind(SemaRef, ExistingClauses, Clause)) + + if (DisallowSinceLastDeviceType(Clause)) return nullptr; ExprResult LoopCount = SemaRef.CheckCollapseLoopCount(Clause.getLoopCount()); @@ -2103,8 +1633,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitCollapseClause( OpenACCClause *SemaOpenACCClauseVisitor::VisitBindClause( SemaOpenACC::OpenACCParsedClause &Clause) { - if (checkAlreadyHasClauseOfKind(SemaRef, ExistingClauses, Clause)) - return nullptr; if (std::holds_alternative(Clause.getBindDetails())) return OpenACCBindClause::Create( @@ -2182,13 +1710,16 @@ SemaOpenACC::ActOnClause(ArrayRef ExistingClauses, if (Clause.getClauseKind() == OpenACCClauseKind::Invalid) return nullptr; - // Diagnose that we don't support this clause on this directive. - if (!doesClauseApplyToDirective(Clause.getDirectiveKind(), - Clause.getClauseKind())) { - Diag(Clause.getBeginLoc(), diag::err_acc_clause_appertainment) - << Clause.getDirectiveKind() << Clause.getClauseKind(); + if (DiagnoseAllowedClauses(Clause.getDirectiveKind(), Clause.getClauseKind(), + Clause.getBeginLoc())) return nullptr; - } + //// Diagnose that we don't support this clause on this directive. + // if (!doesClauseApplyToDirective(Clause.getDirectiveKind(), + // Clause.getClauseKind())) { + // Diag(Clause.getBeginLoc(), diag::err_acc_clause_appertainment) + // << Clause.getDirectiveKind() << Clause.getClauseKind(); + // return nullptr; + // } if (const auto *DevTypeClause = llvm::find_if( ExistingClauses, llvm::IsaPred); diff --git a/clang/lib/Sema/SemaOpenACCClauseAppertainment.cpp b/clang/lib/Sema/SemaOpenACCClauseAppertainment.cpp new file mode 100644 index 0000000000000..6069c77d01676 --- /dev/null +++ b/clang/lib/Sema/SemaOpenACCClauseAppertainment.cpp @@ -0,0 +1,229 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +/// \file +/// This file implements conversions from the ACC.td from the backend to +/// determine appertainment, required/etc. +/// +//===----------------------------------------------------------------------===// + +#include "clang/Basic/DiagnosticSema.h" +#include "clang/Sema/SemaOpenACC.h" + +#include "llvm/ADT/STLExtras.h" +#include "llvm/ADT/bit.h" + +using namespace clang; + +namespace { +// Implements a simple 'enum-set' which stores enum values in a single 64 bit +// value. Flang has `EnumSet` which is pretty sizable/has a lot of dependencies, +// so likely not worth bringing in for this use. +class AccClauseSet { + // We're just using a uint64_t as our underlying rep, so if this size ever + // gets bigger than 64, we probably need a pair of uint64_ts. + static_assert(static_cast(OpenACCClauseKind::Invalid) < 64); + uint64_t Data; + + void setBit(OpenACCClauseKind C) { + Data |= static_cast(1) << static_cast(C); + } + +public: + constexpr AccClauseSet( + const std::initializer_list &Clauses) + : Data(0) { + for (OpenACCClauseKind C : Clauses) + setBit(C); + } + + constexpr bool isSet(OpenACCClauseKind C) const { + return ((Data >> static_cast(C)) & 1) != 0; + } + + void clearBit(OpenACCClauseKind C) { + Data &= ~(static_cast(1) << static_cast(C)); + } + + constexpr bool isEmpty() const { return Data == 0; } + + unsigned popcount() const { return llvm::popcount(Data); } +}; + +struct LLVMClauseLists { + AccClauseSet Allowed; + AccClauseSet AllowedOnce; + AccClauseSet AllowedExclusive; + AccClauseSet Required; +}; +struct LLVMDirectiveClauseRelationships { + OpenACCDirectiveKind DirKind; + LLVMClauseLists Lists; +}; + +} // namespace + +// This introduces these in a llvm::acc namespace, so make sure this stays in +// the global namespace. +#define GEN_CLANG_DIRECTIVE_CLAUSE_SETS +#include "llvm/Frontend/OpenACC/ACC.inc" + +namespace { +LLVMDirectiveClauseRelationships Relations[] = +#define GEN_CLANG_DIRECTIVE_CLAUSE_MAP +#include "llvm/Frontend/OpenACC/ACC.inc" + ; + +const LLVMClauseLists &getListsForDirective(OpenACCDirectiveKind DK) { + + auto Res = llvm::find_if(Relations, + [=](const LLVMDirectiveClauseRelationships &Rel) { + return Rel.DirKind == DK; + }); + assert(Res != std::end(Relations) && "Unknown directive kind?"); + + return Res->Lists; +} + +std::string getListOfClauses(AccClauseSet Set) { + // We could probably come up with a better way to do this smuggling, but this + // is good enough for now. + std::string Output; + llvm::raw_string_ostream OS{Output}; + + for (unsigned I = 0; I < static_cast(OpenACCClauseKind::Invalid); + ++I) { + OpenACCClauseKind CurClause = static_cast(I); + if (!Set.isSet(CurClause)) + continue; + + OS << '\'' << CurClause << '\''; + + Set.clearBit(CurClause); + + if (Set.isEmpty()) { + OS.flush(); + return OS.str(); + } + + OS << ", "; + + if (Set.popcount() == 1) + OS << "or "; + } + OS.flush(); + return OS.str(); +} + +OpenACCClauseKind dealiasClauseKind(OpenACCClauseKind CK) { + switch (CK) { + default: + return CK; +#define VISIT_CLAUSE(NAME) +#define CLAUSE_ALIAS(ALIAS, NAME, DEPRECATED) \ + case OpenACCClauseKind::ALIAS: \ + return OpenACCClauseKind::NAME; +#include "clang/Basic/OpenACCClauses.def" + } + + return CK; +} +} // namespace + +// Diagnoses if `Clauses` list doesn't have at least one of the required +// clauses. +bool SemaOpenACC::DiagnoseRequiredClauses( + OpenACCDirectiveKind DK, SourceLocation DirectiveLoc, + ArrayRef Clauses) { + if (DK == OpenACCDirectiveKind::Invalid) + return false; + + const LLVMClauseLists &Lists = getListsForDirective(DK); + + if (Lists.Required.isEmpty()) + return false; + + for (auto *C : Clauses) { + if (Lists.Required.isSet(dealiasClauseKind(C->getClauseKind()))) + return false; + } + + return Diag(DirectiveLoc, diag::err_acc_construct_one_clause_of) + << DK << getListOfClauses(Lists.Required); + return true; +} + +// Diagnoses a 'CK' on a 'DK' present more than once in a clause-list when it +// isn't allowed. +bool SemaOpenACC::DiagnoseAllowedOnceClauses( + OpenACCDirectiveKind DK, OpenACCClauseKind CK, SourceLocation ClauseLoc, + ArrayRef Clauses) { + if (DK == OpenACCDirectiveKind::Invalid || CK == OpenACCClauseKind::Invalid) + return false; + + OpenACCClauseKind Dealiased = dealiasClauseKind(CK); + + const LLVMClauseLists &Lists = getListsForDirective(DK); + if (!Lists.AllowedOnce.isSet(CK)) + return false; + + auto Res = llvm::find_if(Clauses, [=](const OpenACCClause *C) { + return dealiasClauseKind(C->getClauseKind()) == Dealiased; + }); + + if (Res == Clauses.end()) + return false; + + Diag(ClauseLoc, diag::err_acc_duplicate_clause_disallowed) << DK << CK; + Diag((*Res)->getBeginLoc(), diag::note_acc_previous_clause_here); + return true; +} + +// Diagnoses a 'CK' on a 'DK' being added that isn't allowed to, because another +// clause in 'Clauses' already exists. +bool SemaOpenACC::DiagnoseExclusiveClauses( + OpenACCDirectiveKind DK, OpenACCClauseKind CK, SourceLocation ClauseLoc, + ArrayRef Clauses) { + if (DK == OpenACCDirectiveKind::Invalid || CK == OpenACCClauseKind::Invalid) + return false; + + const LLVMClauseLists &Lists = getListsForDirective(DK); + OpenACCClauseKind Dealiased = dealiasClauseKind(CK); + + // If this isn't on the list, this is fine. + if (!Lists.AllowedExclusive.isSet(Dealiased)) + return false; + + for (const OpenACCClause *C : Clauses) { + if (Lists.AllowedExclusive.isSet(dealiasClauseKind(C->getClauseKind()))) { + Diag(ClauseLoc, diag::err_acc_clause_cannot_combine) + << CK << C->getClauseKind() << DK; + Diag(C->getBeginLoc(), diag::note_acc_previous_clause_here); + + return true; + } + } + + return false; +} + +// Diagnoses if 'CK' is not allowed on a directive of 'DK'. +bool SemaOpenACC::DiagnoseAllowedClauses(OpenACCDirectiveKind DK, + OpenACCClauseKind CK, + SourceLocation ClauseLoc) { + if (DK == OpenACCDirectiveKind::Invalid || CK == OpenACCClauseKind::Invalid) + return false; + const LLVMClauseLists &Lists = getListsForDirective(DK); + OpenACCClauseKind Dealiased = dealiasClauseKind(CK); + + if (!Lists.Allowed.isSet(Dealiased) && !Lists.AllowedOnce.isSet(Dealiased) && + !Lists.AllowedExclusive.isSet(Dealiased) && + !Lists.Required.isSet(Dealiased)) + return Diag(ClauseLoc, diag::err_acc_clause_appertainment) << DK << CK; + + return false; +} diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h index 2469991bf2ce8..b9a09ade37c12 100644 --- a/clang/lib/Sema/TreeTransform.h +++ b/clang/lib/Sema/TreeTransform.h @@ -12355,9 +12355,6 @@ void OpenACCClauseTransform::VisitCollapseClause( NewLoopCount = Self.getSema().OpenACC().CheckCollapseLoopCount(NewLoopCount.get()); - if (!NewLoopCount.isUsable()) - return; - ParsedClause.setCollapseDetails(C.hasForce(), NewLoopCount.get()); NewClause = OpenACCCollapseClause::Create( Self.getSema().getASTContext(), ParsedClause.getBeginLoc(), diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp index d683b7930caf5..70b54b7296882 100644 --- a/clang/lib/Serialization/ASTReader.cpp +++ b/clang/lib/Serialization/ASTReader.cpp @@ -12906,6 +12906,7 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() { return OpenACCBindClause::Create(getContext(), BeginLoc, LParenLoc, readIdentifier(), EndLoc); } + case OpenACCClauseKind::Shortloop: case OpenACCClauseKind::Invalid: llvm_unreachable("Clause serialization not yet implemented"); } diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp index e969214e633f3..ffbe2c4e4d0f1 100644 --- a/clang/lib/Serialization/ASTWriter.cpp +++ b/clang/lib/Serialization/ASTWriter.cpp @@ -8871,6 +8871,7 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) { return; } case OpenACCClauseKind::Invalid: + case OpenACCClauseKind::Shortloop: llvm_unreachable("Clause serialization not yet implemented"); } llvm_unreachable("Invalid Clause Kind"); diff --git a/clang/test/AST/ast-print-openacc-set-construct.cpp b/clang/test/AST/ast-print-openacc-set-construct.cpp index 869a3feb2dd93..1ad6adb17e360 100644 --- a/clang/test/AST/ast-print-openacc-set-construct.cpp +++ b/clang/test/AST/ast-print-openacc-set-construct.cpp @@ -15,8 +15,8 @@ void uses() { #pragma acc set if (Int == 5) device_type(multicore) device_num(Int) // CHECK: #pragma acc set default_async(Int) #pragma acc set default_async(Int) -// CHECK: #pragma acc set if(Int == 5) -#pragma acc set if (Int == 5) +// CHECK: #pragma acc set if(Int == 5) device_type(multicore) +#pragma acc set if (Int == 5) device_type(multicore) // CHECK: #pragma acc set device_type(radeon) #pragma acc set device_type(radeon) // CHECK: #pragma acc set device_num(Int) diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c index ef0bccf390297..716d9cffab3f6 100644 --- a/clang/test/ParserOpenACC/parse-clauses.c +++ b/clang/test/ParserOpenACC/parse-clauses.c @@ -4,21 +4,21 @@ void func() { - // expected-error@+1{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete' or 'detach' clause}} + // expected-error@+1{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete', or 'detach' clause}} #pragma acc exit data finalize - // expected-error@+1{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete' or 'detach' clause}} + // expected-error@+1{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete', or 'detach' clause}} #pragma acc exit data finalize finalize - // expected-error@+2{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete' or 'detach' clause}} + // expected-error@+2{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete', or 'detach' clause}} // expected-error@+1{{invalid OpenACC clause 'invalid'}} #pragma acc exit data finalize invalid - // expected-error@+2{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete' or 'detach' clause}} + // expected-error@+2{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete', or 'detach' clause}} // expected-error@+1{{invalid OpenACC clause 'invalid'}} #pragma acc exit data finalize invalid invalid finalize - // expected-error@+1{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete' or 'detach' clause}} + // expected-error@+1{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete', or 'detach' clause}} #pragma acc exit data wait finalize // expected-error@+1{{OpenACC 'host_data' construct must have at least one 'use_device' clause}} @@ -27,41 +27,41 @@ void func() { // expected-error@+1{{OpenACC 'host_data' construct must have at least one 'use_device' clause}} #pragma acc host_data if_present, if_present - // expected-error@+4{{OpenACC clause 'independent' on 'loop' construct conflicts with previous data dependence clause}} + // expected-error@+4{{OpenACC clause 'independent' may not appear on the same construct as a 'seq' clause on a 'loop' construct}} // expected-note@+3{{previous clause is here}} - // expected-error@+2{{OpenACC clause 'auto' on 'loop' construct conflicts with previous data dependence clause}} + // expected-error@+2{{OpenACC clause 'auto' may not appear on the same construct as a 'seq' clause on a 'loop' construct}} // expected-note@+1{{previous clause is here}} #pragma acc loop seq independent auto for(int i = 0; i < 5;++i) {} - // expected-error@+4{{OpenACC clause 'independent' on 'loop' construct conflicts with previous data dependence clause}} + // expected-error@+4{{OpenACC clause 'independent' may not appear on the same construct as a 'seq' clause on a 'loop' construct}} // expected-note@+3{{previous clause is here}} - // expected-error@+2{{OpenACC clause 'auto' on 'loop' construct conflicts with previous data dependence clause}} + // expected-error@+2{{OpenACC clause 'auto' may not appear on the same construct as a 'seq' clause on a 'loop' construct}} // expected-note@+1{{previous clause is here}} #pragma acc loop seq, independent auto for(int i = 0; i < 5;++i) {} - // expected-error@+4{{OpenACC clause 'independent' on 'loop' construct conflicts with previous data dependence clause}} + // expected-error@+4{{OpenACC clause 'independent' may not appear on the same construct as a 'seq' clause on a 'loop' construct}} // expected-note@+3{{previous clause is here}} - // expected-error@+2{{OpenACC clause 'auto' on 'loop' construct conflicts with previous data dependence clause}} + // expected-error@+2{{OpenACC clause 'auto' may not appear on the same construct as a 'seq' clause on a 'loop' construct}} // expected-note@+1{{previous clause is here}} #pragma acc loop seq independent, auto for(int i = 0; i < 5;++i) {} - // expected-error@+3{{OpenACC clause 'independent' on 'kernels loop' construct conflicts with previous data dependence clause}} - // expected-error@+2{{OpenACC clause 'auto' on 'kernels loop' construct conflicts with previous data dependence clause}} + // expected-error@+3{{OpenACC clause 'independent' may not appear on the same construct as a 'seq' clause on a 'kernels loop' construct}} + // expected-error@+2{{OpenACC clause 'auto' may not appear on the same construct as a 'seq' clause on a 'kernels loop' construct}} // expected-note@+1 2{{previous clause is here}} #pragma acc kernels loop seq independent auto for(int i = 0; i < 5;++i) {} - // expected-error@+3{{OpenACC clause 'independent' on 'serial loop' construct conflicts with previous data dependence clause}} - // expected-error@+2{{OpenACC clause 'auto' on 'serial loop' construct conflicts with previous data dependence clause}} + // expected-error@+3{{OpenACC clause 'independent' may not appear on the same construct as a 'seq' clause on a 'serial loop' construct}} + // expected-error@+2{{OpenACC clause 'auto' may not appear on the same construct as a 'seq' clause on a 'serial loop' construct}} // expected-note@+1 2{{previous clause is here}} #pragma acc serial loop seq, independent auto for(int i = 0; i < 5;++i) {} - // expected-error@+3{{OpenACC clause 'independent' on 'parallel loop' construct conflicts with previous data dependence clause}} - // expected-error@+2{{OpenACC clause 'auto' on 'parallel loop' construct conflicts with previous data dependence clause}} + // expected-error@+3{{OpenACC clause 'independent' may not appear on the same construct as a 'seq' clause on a 'parallel loop' construct}} + // expected-error@+2{{OpenACC clause 'auto' may not appear on the same construct as a 'seq' clause on a 'parallel loop' construct}} // expected-note@+1 2{{previous clause is here}} #pragma acc parallel loop seq independent, auto for(int i = 0; i < 5;++i) {} @@ -860,20 +860,20 @@ void IntExprParsing() { #pragma acc init device_num(returns_int()) // expected-error@+2{{expected '('}} - // expected-error@+1{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', 'device_type' or 'if' clause}} + // expected-error@+1{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', or 'device_type' clause}} #pragma acc set default_async // expected-error@+2{{expected expression}} - // expected-error@+1{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', 'device_type' or 'if' clause}} + // expected-error@+1{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', or 'device_type' clause}} #pragma acc set default_async() // expected-error@+2{{use of undeclared identifier 'invalid'}} - // expected-error@+1{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', 'device_type' or 'if' clause}} + // expected-error@+1{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', or 'device_type' clause}} #pragma acc set default_async(invalid) // expected-error@+3{{expected ')'}} // expected-note@+2{{to match this '('}} - // expected-error@+1{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', 'device_type' or 'if' clause}} + // expected-error@+1{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', or 'device_type' clause}} #pragma acc set default_async(5, 4) #pragma acc set default_async(5) diff --git a/clang/test/ParserOpenACC/parse-constructs.c b/clang/test/ParserOpenACC/parse-constructs.c index 71e47abd3a825..0365e3ab83d8d 100644 --- a/clang/test/ParserOpenACC/parse-constructs.c +++ b/clang/test/ParserOpenACC/parse-constructs.c @@ -54,15 +54,15 @@ void func() { // expected-error@+1{{invalid OpenACC clause 'clause'}} #pragma acc kernels clause list for(;;){} - // expected-error@+2{{OpenACC 'data' construct must have at least one 'copy', 'copyin', 'copyout', 'create', 'no_create', 'present', 'deviceptr', 'attach' or 'default' clause}} + // expected-error@+2{{OpenACC 'data' construct must have at least one 'attach', 'copy', 'copyin', 'copyout', 'create', 'default', 'deviceptr', 'no_create', or 'present' clause}} // expected-error@+1{{invalid OpenACC clause 'clause'}} #pragma acc data clause list for(;;){} - // expected-error@+2{{OpenACC 'enter data' construct must have at least one 'copyin', 'create' or 'attach' clause}} + // expected-error@+2{{OpenACC 'enter data' construct must have at least one 'attach', 'copyin', or 'create' clause}} // expected-error@+1{{invalid OpenACC clause 'clause'}} #pragma acc enter data clause list for(;;){} - // expected-error@+2{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete' or 'detach' clause}} + // expected-error@+2{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete', or 'detach' clause}} // expected-error@+1{{invalid OpenACC clause 'clause'}} #pragma acc exit data clause list for(;;){} @@ -141,10 +141,10 @@ void func() { #pragma acc shutdown clause list for(;;){} // expected-error@+2{{invalid OpenACC clause 'clause'}} - // expected-error@+1{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', 'device_type' or 'if' clause}} + // expected-error@+1{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', or 'device_type' clause}} #pragma acc set clause list for(;;){} - // expected-error@+2{{OpenACC 'update' construct must have at least one 'self', 'host' or 'device' clause}} + // expected-error@+2{{OpenACC 'update' construct must have at least one 'device', 'host', or 'self' clause}} // expected-error@+1{{invalid OpenACC clause 'clause'}} #pragma acc update clause list for(;;){} @@ -153,7 +153,7 @@ void func() { #pragma acc routine seq void routine_func(); // expected-error@+2{{invalid OpenACC clause 'clause'}} -// expected-error@+1{{OpenACC 'routine' construct must have at least one 'gang', 'worker', 'vector' or 'seq' clause}} +// expected-error@+1{{OpenACC 'routine' construct must have at least one 'gang', 'seq', 'vector', or 'worker' clause}} #pragma acc routine clause list void routine_func(); @@ -161,12 +161,12 @@ void routine_func(); #pragma acc routine (func_name) seq // expected-error@+3{{use of undeclared identifier 'func_name'}} // expected-error@+2{{invalid OpenACC clause 'clause'}} -// expected-error@+1{{OpenACC 'routine' construct must have at least one 'gang', 'worker', 'vector' or 'seq' clause}} +// expected-error@+1{{OpenACC 'routine' construct must have at least one 'gang', 'seq', 'vector', or 'worker' clause}} #pragma acc routine (func_name) clause list #pragma acc routine (routine_func) seq // expected-error@+2{{invalid OpenACC clause 'clause'}} -// expected-error@+1{{OpenACC 'routine' construct must have at least one 'gang', 'worker', 'vector' or 'seq' clause}} +// expected-error@+1{{OpenACC 'routine' construct must have at least one 'gang', 'seq', 'vector', or 'worker' clause}} #pragma acc routine (routine_func) clause list // expected-error@+2{{expected ')'}} diff --git a/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c b/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c index dbba86a1ffc61..5c64c621b0e7e 100644 --- a/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c +++ b/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c @@ -8,27 +8,27 @@ void uses() { #pragma acc parallel loop independent for(unsigned i = 0; i < 5; ++i); - // expected-error@+2{{OpenACC clause 'seq' on 'parallel loop' construct conflicts with previous data dependence clause}} + // expected-error@+2{{OpenACC clause 'seq' may not appear on the same construct as a 'auto' clause on a 'parallel loop' construct}} // expected-note@+1{{previous clause is here}} #pragma acc parallel loop auto seq for(unsigned i = 0; i < 5; ++i); - // expected-error@+2{{OpenACC clause 'independent' on 'parallel loop' construct conflicts with previous data dependence clause}} + // expected-error@+2{{OpenACC clause 'independent' may not appear on the same construct as a 'auto' clause on a 'parallel loop' construct}} // expected-note@+1{{previous clause is here}} #pragma acc parallel loop auto independent for(unsigned i = 0; i < 5; ++i); - // expected-error@+2{{OpenACC clause 'auto' on 'parallel loop' construct conflicts with previous data dependence clause}} + // expected-error@+2{{OpenACC clause 'auto' may not appear on the same construct as a 'seq' clause on a 'parallel loop' construct}} // expected-note@+1{{previous clause is here}} #pragma acc parallel loop seq auto for(unsigned i = 0; i < 5; ++i); - // expected-error@+2{{OpenACC clause 'independent' on 'parallel loop' construct conflicts with previous data dependence clause}} + // expected-error@+2{{OpenACC clause 'independent' may not appear on the same construct as a 'seq' clause on a 'parallel loop' construct}} // expected-note@+1{{previous clause is here}} #pragma acc parallel loop seq independent for(unsigned i = 0; i < 5; ++i); - // expected-error@+2{{OpenACC clause 'auto' on 'parallel loop' construct conflicts with previous data dependence clause}} + // expected-error@+2{{OpenACC clause 'auto' may not appear on the same construct as a 'independent' clause on a 'parallel loop' construct}} // expected-note@+1{{previous clause is here}} #pragma acc parallel loop independent auto for(unsigned i = 0; i < 5; ++i); - // expected-error@+2{{OpenACC clause 'seq' on 'parallel loop' construct conflicts with previous data dependence clause}} + // expected-error@+2{{OpenACC clause 'seq' may not appear on the same construct as a 'independent' clause on a 'parallel loop' construct}} // expected-note@+1{{previous clause is here}} #pragma acc parallel loop independent seq for(unsigned i = 0; i < 5; ++i); diff --git a/clang/test/SemaOpenACC/combined-construct-collapse-clause.cpp b/clang/test/SemaOpenACC/combined-construct-collapse-clause.cpp index f2be6622007c3..3f8fb4203d0e5 100644 --- a/clang/test/SemaOpenACC/combined-construct-collapse-clause.cpp +++ b/clang/test/SemaOpenACC/combined-construct-collapse-clause.cpp @@ -214,7 +214,7 @@ void no_other_directives() { #pragma acc serial loop collapse(2) for(unsigned i = 0; i < 5; ++i) { for(unsigned j = 0; j < 5; ++j) { - // expected-error@+1{{OpenACC 'data' construct must have at least one 'copy', 'copyin', 'copyout', 'create', 'no_create', 'present', 'deviceptr', 'attach' or 'default' clause}} + // expected-error@+1{{OpenACC 'data' construct must have at least one 'attach', 'copy', 'copyin', 'copyout', 'create', 'default', 'deviceptr', 'no_create', or 'present' clause}} #pragma acc data ; } @@ -222,7 +222,7 @@ void no_other_directives() { // expected-note@+1{{active 'collapse' clause defined here}} #pragma acc kernels loop collapse(2) for(unsigned i = 0; i < 5; ++i) { - // expected-error@+2{{OpenACC 'data' construct must have at least one 'copy', 'copyin', 'copyout', 'create', 'no_create', 'present', 'deviceptr', 'attach' or 'default' clause}} + // expected-error@+2{{OpenACC 'data' construct must have at least one 'attach', 'copy', 'copyin', 'copyout', 'create', 'default', 'deviceptr', 'no_create', or 'present' clause}} // expected-error@+1{{OpenACC 'data' construct cannot appear in intervening code of a 'kernels loop' with a 'collapse' clause}} #pragma acc data for(unsigned j = 0; j < 5; ++j) { diff --git a/clang/test/SemaOpenACC/combined-construct-if-clause.c b/clang/test/SemaOpenACC/combined-construct-if-clause.c index 5580fdfe22b7b..d5e42d6105963 100644 --- a/clang/test/SemaOpenACC/combined-construct-if-clause.c +++ b/clang/test/SemaOpenACC/combined-construct-if-clause.c @@ -43,7 +43,7 @@ void BoolExpr(int *I, float *F) { #pragma acc kernels loop if (*I < *F) for (unsigned i = 0; i < 5; ++i); - // expected-error@+1{{OpenACC 'data' construct must have at least one 'copy', 'copyin', 'copyout', 'create', 'no_create', 'present', 'deviceptr', 'attach' or 'default' clause}} + // expected-error@+1{{OpenACC 'data' construct must have at least one 'attach', 'copy', 'copyin', 'copyout', 'create', 'default', 'deviceptr', 'no_create', or 'present' clause}} #pragma acc data if (*I < *F) for (unsigned i = 0; i < 5; ++i); #pragma acc parallel loop if (*I < *F) diff --git a/clang/test/SemaOpenACC/combined-construct-num_gangs-clause.c b/clang/test/SemaOpenACC/combined-construct-num_gangs-clause.c index bd035bd4a51a2..cd06f5e05ee69 100644 --- a/clang/test/SemaOpenACC/combined-construct-num_gangs-clause.c +++ b/clang/test/SemaOpenACC/combined-construct-num_gangs-clause.c @@ -13,6 +13,31 @@ void Test() { #pragma acc parallel loop num_gangs(1) for(int i = 5; i < 10;++i); + // expected-error@+2{{OpenACC 'num_gangs' clause cannot appear more than once on a 'kernels loop' directive}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels loop num_gangs(1) num_gangs(2) + for(int i = 5; i < 10;++i); + + // expected-error@+2{{OpenACC 'num_gangs' clause cannot appear more than once on a 'parallel loop' directive}} + // expected-note@+1{{previous clause is here}} +#pragma acc parallel loop num_gangs(1) num_gangs(2) + for(int i = 5; i < 10;++i); + + // expected-error@+3{{OpenACC 'num_gangs' clause cannot appear more than once in a 'device_type' region on a 'kernels loop' directive}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels loop num_gangs(1) device_type(*) num_gangs(1) num_gangs(2) + for(int i = 5; i < 10;++i); + + // expected-error@+3{{OpenACC 'num_gangs' clause cannot appear more than once in a 'device_type' region on a 'parallel loop' directive}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc parallel loop device_type(*) num_gangs(1) num_gangs(2) + for(int i = 5; i < 10;++i); + +#pragma acc parallel loop num_gangs(1) device_type(*) num_gangs(2) + for(int i = 5; i < 10;++i); + // expected-error@+1{{OpenACC clause 'num_gangs' requires expression of integer type}} #pragma acc parallel loop num_gangs(getF()) for(int i = 5; i < 10;++i); @@ -25,16 +50,6 @@ void Test() { #pragma acc parallel loop num_gangs() for(int i = 5; i < 10;++i); - // expected-error@+2{{OpenACC 'num_gangs' clause cannot appear more than once on a 'kernels loop' directive}} - // expected-note@+1{{previous clause is here}} -#pragma acc kernels loop num_gangs(1) num_gangs(2) - for(int i = 5; i < 10;++i); - - // expected-error@+2{{OpenACC 'num_gangs' clause cannot appear more than once on a 'parallel loop' directive}} - // expected-note@+1{{previous clause is here}} -#pragma acc parallel loop num_gangs(1) num_gangs(2) - for(int i = 5; i < 10;++i); - // expected-error@+1{{too many integer expression arguments provided to OpenACC 'num_gangs' clause: 'kernels loop' directive expects maximum of 1, 2 were provided}} #pragma acc kernels loop num_gangs(1, getS()) for(int i = 5; i < 10;++i); diff --git a/clang/test/SemaOpenACC/combined-construct-num_workers-clause.c b/clang/test/SemaOpenACC/combined-construct-num_workers-clause.c index a5891f071bb03..c3b648e438081 100644 --- a/clang/test/SemaOpenACC/combined-construct-num_workers-clause.c +++ b/clang/test/SemaOpenACC/combined-construct-num_workers-clause.c @@ -13,6 +13,31 @@ void Test() { #pragma acc parallel loop num_workers(1) for(int i = 5; i < 10;++i); + // expected-error@+2{{OpenACC 'num_workers' clause cannot appear more than once on a 'kernels loop' directive}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels loop num_workers(1) num_workers(2) + for(int i = 5; i < 10;++i); + + // expected-error@+2{{OpenACC 'num_workers' clause cannot appear more than once on a 'parallel loop' directive}} + // expected-note@+1{{previous clause is here}} +#pragma acc parallel loop num_workers(1) num_workers(2) + for(int i = 5; i < 10;++i); + + // expected-error@+3{{OpenACC 'num_workers' clause cannot appear more than once in a 'device_type' region on a 'kernels loop' directive}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels loop num_workers(1) device_type(*) num_workers(1) num_workers(2) + for(int i = 5; i < 10;++i); + + // expected-error@+3{{OpenACC 'num_workers' clause cannot appear more than once in a 'device_type' region on a 'parallel loop' directive}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc parallel loop device_type(*) num_workers(1) num_workers(2) + for(int i = 5; i < 10;++i); + +#pragma acc parallel loop num_workers(1) device_type(*) num_workers(2) + for(int i = 5; i < 10;++i); + // expected-error@+1{{OpenACC clause 'num_workers' requires expression of integer type}} #pragma acc parallel loop num_workers(getF()) for(int i = 5; i < 10;++i); diff --git a/clang/test/SemaOpenACC/combined-construct-tile-clause.cpp b/clang/test/SemaOpenACC/combined-construct-tile-clause.cpp index 00c551e163666..deab634c39886 100644 --- a/clang/test/SemaOpenACC/combined-construct-tile-clause.cpp +++ b/clang/test/SemaOpenACC/combined-construct-tile-clause.cpp @@ -143,13 +143,6 @@ void only_for_loops() { do{}while(true); } -void only_one_on_loop() { - // expected-error@+2{{OpenACC 'tile' clause cannot appear more than once on a 'serial loop' directive}} - // expected-note@+1{{previous clause is here}} -#pragma acc serial loop tile(1) tile(1) - for(int i = 0; i < 5; ++i); -} - template void depth_too_high_templ() { // expected-error@+1{{'tile' clause specifies a loop count greater than the number of available loops}} diff --git a/clang/test/SemaOpenACC/combined-construct-vector_length-clause.c b/clang/test/SemaOpenACC/combined-construct-vector_length-clause.c index 8b6dedd9b83ba..72ab0126b1113 100644 --- a/clang/test/SemaOpenACC/combined-construct-vector_length-clause.c +++ b/clang/test/SemaOpenACC/combined-construct-vector_length-clause.c @@ -13,6 +13,31 @@ void Test() { #pragma acc parallel loop vector_length(1) for(int i = 5; i < 10;++i); + // expected-error@+2{{OpenACC 'vector_length' clause cannot appear more than once on a 'kernels loop' directive}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels loop vector_length(1) vector_length(2) + for(int i = 5; i < 10;++i); + + // expected-error@+2{{OpenACC 'vector_length' clause cannot appear more than once on a 'parallel loop' directive}} + // expected-note@+1{{previous clause is here}} +#pragma acc parallel loop vector_length(1) vector_length(2) + for(int i = 5; i < 10;++i); + + // expected-error@+3{{OpenACC 'vector_length' clause cannot appear more than once in a 'device_type' region on a 'kernels loop' directive}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels loop vector_length(1) device_type(*) vector_length(1) vector_length(2) + for(int i = 5; i < 10;++i); + + // expected-error@+3{{OpenACC 'vector_length' clause cannot appear more than once in a 'device_type' region on a 'parallel loop' directive}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc parallel loop device_type(*) vector_length(1) vector_length(2) + for(int i = 5; i < 10;++i); + +#pragma acc parallel loop vector_length(1) device_type(*) vector_length(2) + for(int i = 5; i < 10;++i); + // expected-error@+1{{OpenACC clause 'vector_length' requires expression of integer type}} #pragma acc parallel loop vector_length(getF()) for(int i = 5; i < 10;++i); diff --git a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c index 16caa7205721d..2ff6369383714 100644 --- a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c +++ b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c @@ -34,10 +34,10 @@ void uses() { #pragma acc kernels dtype(MACRO) while(1); - // expected-error@+2{{OpenACC 'enter data' construct must have at least one 'copyin', 'create' or 'attach' clause}} + // expected-error@+2{{OpenACC 'enter data' construct must have at least one 'attach', 'copyin', or 'create' clause}} // expected-error@+1{{OpenACC 'device_type' clause is not valid on 'enter data' directive}} #pragma acc enter data device_type(I) - // expected-error@+2{{OpenACC 'enter data' construct must have at least one 'copyin', 'create' or 'attach' clause}} + // expected-error@+2{{OpenACC 'enter data' construct must have at least one 'attach', 'copyin', or 'create' clause}} // expected-error@+1{{OpenACC 'dtype' clause is not valid on 'enter data' directive}} #pragma acc enter data dtype(I) diff --git a/clang/test/SemaOpenACC/compute-construct-if-clause.c b/clang/test/SemaOpenACC/compute-construct-if-clause.c index c0ea88f06284d..0064303e1e217 100644 --- a/clang/test/SemaOpenACC/compute-construct-if-clause.c +++ b/clang/test/SemaOpenACC/compute-construct-if-clause.c @@ -43,7 +43,7 @@ void BoolExpr(int *I, float *F) { #pragma acc kernels if (*I < *F) while(0); - // expected-error@+1{{OpenACC 'data' construct must have at least one 'copy', 'copyin', 'copyout', 'create', 'no_create', 'present', 'deviceptr', 'attach' or 'default' clause}} + // expected-error@+1{{OpenACC 'data' construct must have at least one 'attach', 'copy', 'copyin', 'copyout', 'create', 'default', 'deviceptr', 'no_create', or 'present' clause}} #pragma acc data if (*I < *F) while(0); #pragma acc parallel loop if (*I < *F) diff --git a/clang/test/SemaOpenACC/compute-construct-num_gangs-clause.c b/clang/test/SemaOpenACC/compute-construct-num_gangs-clause.c index 40b3b652d21da..9651514ca0fff 100644 --- a/clang/test/SemaOpenACC/compute-construct-num_gangs-clause.c +++ b/clang/test/SemaOpenACC/compute-construct-num_gangs-clause.c @@ -22,6 +22,21 @@ void Test() { #pragma acc parallel num_gangs(1) num_gangs(2) while(1); + // expected-error@+3{{OpenACC 'num_gangs' clause cannot appear more than once in a 'device_type' region on a 'kernels' directive}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels num_gangs(1) device_type(*) num_gangs(1) num_gangs(2) + while(1); + + // expected-error@+3{{OpenACC 'num_gangs' clause cannot appear more than once in a 'device_type' region on a 'parallel' directive}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc parallel device_type(*) num_gangs(1) num_gangs(2) + while(1); + +#pragma acc parallel num_gangs(1) device_type(*) num_gangs(2) + while(1); + // expected-error@+1{{too many integer expression arguments provided to OpenACC 'num_gangs' clause: 'kernels' directive expects maximum of 1, 2 were provided}} #pragma acc kernels num_gangs(1, getS()) while(1); diff --git a/clang/test/SemaOpenACC/compute-construct-num_workers-clause.c b/clang/test/SemaOpenACC/compute-construct-num_workers-clause.c index a78e2c4feadbb..23a73f0bf49bd 100644 --- a/clang/test/SemaOpenACC/compute-construct-num_workers-clause.c +++ b/clang/test/SemaOpenACC/compute-construct-num_workers-clause.c @@ -8,6 +8,31 @@ void Test() { #pragma acc kernels num_workers(1) while(1); + // expected-error@+2{{OpenACC 'num_workers' clause cannot appear more than once on a 'kernels' directive}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels num_workers(1) num_workers(2) + while(1); + + // expected-error@+2{{OpenACC 'num_workers' clause cannot appear more than once on a 'parallel' directive}} + // expected-note@+1{{previous clause is here}} +#pragma acc parallel num_workers(1) num_workers(2) + while(1); + + // expected-error@+3{{OpenACC 'num_workers' clause cannot appear more than once in a 'device_type' region on a 'kernels' directive}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels num_workers(1) device_type(*) num_workers(1) num_workers(2) + while(1); + + // expected-error@+3{{OpenACC 'num_workers' clause cannot appear more than once in a 'device_type' region on a 'parallel' directive}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc parallel device_type(*) num_workers(1) num_workers(2) + while(1); + +#pragma acc parallel num_workers(1) device_type(*) num_workers(2) + while(1); + // expected-error@+1{{OpenACC 'num_workers' clause is not valid on 'serial' directive}} #pragma acc serial num_workers(1) while(1); diff --git a/clang/test/SemaOpenACC/compute-construct-vector_length-clause.c b/clang/test/SemaOpenACC/compute-construct-vector_length-clause.c index c9083df0df722..c79fc627252f1 100644 --- a/clang/test/SemaOpenACC/compute-construct-vector_length-clause.c +++ b/clang/test/SemaOpenACC/compute-construct-vector_length-clause.c @@ -8,6 +8,31 @@ void Test() { #pragma acc kernels vector_length(1) while(1); + // expected-error@+2{{OpenACC 'vector_length' clause cannot appear more than once on a 'kernels' directive}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels vector_length(1) vector_length(2) + while(1); + + // expected-error@+2{{OpenACC 'vector_length' clause cannot appear more than once on a 'parallel' directive}} + // expected-note@+1{{previous clause is here}} +#pragma acc parallel vector_length(1) vector_length(2) + while(1); + + // expected-error@+3{{OpenACC 'vector_length' clause cannot appear more than once in a 'device_type' region on a 'kernels' directive}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels vector_length(1) device_type(*) vector_length(1) vector_length(2) + while(1); + + // expected-error@+3{{OpenACC 'vector_length' clause cannot appear more than once in a 'device_type' region on a 'parallel' directive}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc parallel device_type(*) vector_length(1) vector_length(2) + while(1); + +#pragma acc parallel vector_length(1) device_type(*) vector_length(2) + while(1); + // expected-error@+1{{OpenACC 'vector_length' clause is not valid on 'serial' directive}} #pragma acc serial vector_length(1) while(1); diff --git a/clang/test/SemaOpenACC/data-construct-copy-clause.c b/clang/test/SemaOpenACC/data-construct-copy-clause.c index 0b2b0534073ed..b19f3f465caaf 100644 --- a/clang/test/SemaOpenACC/data-construct-copy-clause.c +++ b/clang/test/SemaOpenACC/data-construct-copy-clause.c @@ -56,10 +56,10 @@ void uses(int IntParam, short *PointerParam, float ArrayParam[5], Complete Compo #pragma acc data copy((float)ArrayParam[2]) ; - // expected-error@+2{{OpenACC 'enter data' construct must have at least one 'copyin', 'create' or 'attach' clause}} + // expected-error@+2{{OpenACC 'enter data' construct must have at least one 'attach', 'copyin', or 'create' clause}} // expected-error@+1{{OpenACC 'copy' clause is not valid on 'enter data' directive}} #pragma acc enter data copy(LocalInt) - // expected-error@+2{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete' or 'detach' clause}} + // expected-error@+2{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete', or 'detach' clause}} // expected-error@+1{{OpenACC 'pcopy' clause is not valid on 'exit data' directive}} #pragma acc exit data pcopy(LocalInt) // expected-error@+2{{OpenACC 'host_data' construct must have at least one 'use_device' clause}} diff --git a/clang/test/SemaOpenACC/data-construct-copyin-clause.c b/clang/test/SemaOpenACC/data-construct-copyin-clause.c index edc3f0a2e91fe..9b51a28d51a3d 100644 --- a/clang/test/SemaOpenACC/data-construct-copyin-clause.c +++ b/clang/test/SemaOpenACC/data-construct-copyin-clause.c @@ -63,7 +63,7 @@ void uses(int IntParam, short *PointerParam, float ArrayParam[5], Complete Compo #pragma acc data copyin(invalid:(float)ArrayParam[2]) ; - // expected-error@+2{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete' or 'detach' clause}} + // expected-error@+2{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete', or 'detach' clause}} // expected-error@+1{{OpenACC 'copyin' clause is not valid on 'exit data' directive}} #pragma acc exit data copyin(LocalInt) // expected-error@+2{{OpenACC 'host_data' construct must have at least one 'use_device' clause}} diff --git a/clang/test/SemaOpenACC/data-construct-copyout-clause.c b/clang/test/SemaOpenACC/data-construct-copyout-clause.c index 8d137e093db0e..0c2264c56fcb8 100644 --- a/clang/test/SemaOpenACC/data-construct-copyout-clause.c +++ b/clang/test/SemaOpenACC/data-construct-copyout-clause.c @@ -63,7 +63,7 @@ void uses(int IntParam, short *PointerParam, float ArrayParam[5], Complete Compo #pragma acc data copyout(invalid:(float)ArrayParam[2]) ; - // expected-error@+2{{OpenACC 'enter data' construct must have at least one 'copyin', 'create' or 'attach' clause}} + // expected-error@+2{{OpenACC 'enter data' construct must have at least one 'attach', 'copyin', or 'create' clause}} // expected-error@+1{{OpenACC 'copyout' clause is not valid on 'enter data' directive}} #pragma acc enter data copyout(LocalInt) // expected-error@+2{{OpenACC 'host_data' construct must have at least one 'use_device' clause}} diff --git a/clang/test/SemaOpenACC/data-construct-create-clause.c b/clang/test/SemaOpenACC/data-construct-create-clause.c index e49d53b17ee82..560c6b65cc502 100644 --- a/clang/test/SemaOpenACC/data-construct-create-clause.c +++ b/clang/test/SemaOpenACC/data-construct-create-clause.c @@ -63,7 +63,7 @@ void uses(int IntParam, short *PointerParam, float ArrayParam[5], Complete Compo #pragma acc data create(invalid:(float)ArrayParam[2]) ; - // expected-error@+2{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete' or 'detach' clause}} + // expected-error@+2{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete', or 'detach' clause}} // expected-error@+1{{OpenACC 'create' clause is not valid on 'exit data' directive}} #pragma acc exit data create(LocalInt) // expected-error@+2{{OpenACC 'host_data' construct must have at least one 'use_device' clause}} diff --git a/clang/test/SemaOpenACC/data-construct-default-clause.c b/clang/test/SemaOpenACC/data-construct-default-clause.c index 150d3ec22dcda..9e8ee529e9b8a 100644 --- a/clang/test/SemaOpenACC/data-construct-default-clause.c +++ b/clang/test/SemaOpenACC/data-construct-default-clause.c @@ -1,7 +1,7 @@ // RUN: %clang_cc1 %s -fopenacc -verify void use() { - // expected-error@+2{{OpenACC 'data' construct must have at least one 'copy', 'copyin', 'copyout', 'create', 'no_create', 'present', 'deviceptr', 'attach' or 'default' clause}} + // expected-error@+2{{OpenACC 'data' construct must have at least one 'attach', 'copy', 'copyin', 'copyout', 'create', 'default', 'deviceptr', 'no_create', or 'present' clause}} // expected-error@+1{{invalid value for 'default' clause; expected 'present' or 'none'}} #pragma acc data default(garbage) ; @@ -13,11 +13,11 @@ void use() { // expected-note@+1{{previous clause is here}} #pragma acc data default(none) default(present) ; - // expected-error@+2{{OpenACC 'enter data' construct must have at least one 'copyin', 'create' or 'attach' clause}} + // expected-error@+2{{OpenACC 'enter data' construct must have at least one 'attach', 'copyin', or 'create' clause}} // expected-error@+1{{OpenACC 'default' clause is not valid on 'enter data' directive}} #pragma acc enter data default(present) ; - // expected-error@+2{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete' or 'detach' clause}} + // expected-error@+2{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete', or 'detach' clause}} // expected-error@+1{{OpenACC 'default' clause is not valid on 'exit data' directive}} #pragma acc exit data default(none) ; diff --git a/clang/test/SemaOpenACC/data-construct-delete-clause.c b/clang/test/SemaOpenACC/data-construct-delete-clause.c index 05093dbc4e01b..d2727de8dba63 100644 --- a/clang/test/SemaOpenACC/data-construct-delete-clause.c +++ b/clang/test/SemaOpenACC/data-construct-delete-clause.c @@ -37,11 +37,11 @@ void uses(int IntParam, short *PointerParam, float ArrayParam[5], Complete Compo // expected-error@+1{{OpenACC variable is not a valid variable name, sub-array, array element, member of a composite variable, or composite variable member}} #pragma acc exit data delete((float)ArrayParam[2]) - // expected-error@+2{{OpenACC 'data' construct must have at least one 'copy', 'copyin', 'copyout', 'create', 'no_create', 'present', 'deviceptr', 'attach' or 'default' clause}} + // expected-error@+2{{OpenACC 'data' construct must have at least one 'attach', 'copy', 'copyin', 'copyout', 'create', 'default', 'deviceptr', 'no_create', or 'present' clause}} // expected-error@+1{{OpenACC 'delete' clause is not valid on 'data' directive}} #pragma acc data delete(LocalInt) ; - // expected-error@+2{{OpenACC 'enter data' construct must have at least one 'copyin', 'create' or 'attach' clause}} + // expected-error@+2{{OpenACC 'enter data' construct must have at least one 'attach', 'copyin', or 'create' clause}} // expected-error@+1{{OpenACC 'delete' clause is not valid on 'enter data' directive}} #pragma acc enter data delete(LocalInt) // expected-error@+2{{OpenACC 'host_data' construct must have at least one 'use_device' clause}} diff --git a/clang/test/SemaOpenACC/data-construct-device_type-clause.c b/clang/test/SemaOpenACC/data-construct-device_type-clause.c index e450e320fa53c..a8234224540ba 100644 --- a/clang/test/SemaOpenACC/data-construct-device_type-clause.c +++ b/clang/test/SemaOpenACC/data-construct-device_type-clause.c @@ -47,7 +47,7 @@ void uses() { // expected-note@+1{{previous clause is here}} #pragma acc data default(none) device_type(radeon) attach(Var) ; - // expected-error@+3{{OpenACC 'data' construct must have at least one 'copy', 'copyin', 'copyout', 'create', 'no_create', 'present', 'deviceptr', 'attach' or 'default' clause}} + // expected-error@+3{{OpenACC 'data' construct must have at least one 'attach', 'copy', 'copyin', 'copyout', 'create', 'default', 'deviceptr', 'no_create', or 'present' clause}} // expected-error@+2{{OpenACC clause 'default' may not follow a 'device_type' clause in a 'data' construct}} // expected-note@+1{{previous clause is here}} #pragma acc data device_type(radeon) default(none) diff --git a/clang/test/SemaOpenACC/data-construct-no_create-clause.c b/clang/test/SemaOpenACC/data-construct-no_create-clause.c index 0eb459eb0009a..1de39955ba8ec 100644 --- a/clang/test/SemaOpenACC/data-construct-no_create-clause.c +++ b/clang/test/SemaOpenACC/data-construct-no_create-clause.c @@ -48,10 +48,10 @@ void uses(int IntParam, short *PointerParam, float ArrayParam[5], Complete Compo #pragma acc data no_create((float)ArrayParam[2]) ; - // expected-error@+2{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete' or 'detach' clause}} + // expected-error@+2{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete', or 'detach' clause}} // expected-error@+1{{OpenACC 'no_create' clause is not valid on 'exit data' directive}} #pragma acc exit data no_create(LocalInt) - // expected-error@+2{{OpenACC 'enter data' construct must have at least one 'copyin', 'create' or 'attach' clause}} + // expected-error@+2{{OpenACC 'enter data' construct must have at least one 'attach', 'copyin', or 'create' clause}} // expected-error@+1{{OpenACC 'no_create' clause is not valid on 'enter data' directive}} #pragma acc enter data no_create(LocalInt) // expected-error@+2{{OpenACC 'host_data' construct must have at least one 'use_device' clause}} diff --git a/clang/test/SemaOpenACC/data-construct-use_device-clause.c b/clang/test/SemaOpenACC/data-construct-use_device-clause.c index c2e7fd17f8c8d..9239757b8e85d 100644 --- a/clang/test/SemaOpenACC/data-construct-use_device-clause.c +++ b/clang/test/SemaOpenACC/data-construct-use_device-clause.c @@ -52,14 +52,14 @@ void uses(int IntParam, short *PointerParam, float ArrayParam[5], Complete Compo #pragma acc host_data use_device((float)ArrayParam[2]) ; - // expected-error@+2{{OpenACC 'data' construct must have at least one 'copy', 'copyin', 'copyout', 'create', 'no_create', 'present', 'deviceptr', 'attach' or 'default' clause}} + // expected-error@+2{{OpenACC 'data' construct must have at least one 'attach', 'copy', 'copyin', 'copyout', 'create', 'default', 'deviceptr', 'no_create', or 'present' clause}} // expected-error@+1{{OpenACC 'use_device' clause is not valid on 'data' directive}} #pragma acc data use_device(LocalInt) ; - // expected-error@+2{{OpenACC 'enter data' construct must have at least one 'copyin', 'create' or 'attach' clause}} + // expected-error@+2{{OpenACC 'enter data' construct must have at least one 'attach', 'copyin', or 'create' clause}} // expected-error@+1{{OpenACC 'use_device' clause is not valid on 'enter data' directive}} #pragma acc enter data use_device(LocalInt) - // expected-error@+2{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete' or 'detach' clause}} + // expected-error@+2{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete', or 'detach' clause}} // expected-error@+1{{OpenACC 'use_device' clause is not valid on 'exit data' directive}} #pragma acc exit data use_device(LocalInt) } diff --git a/clang/test/SemaOpenACC/data-construct.cpp b/clang/test/SemaOpenACC/data-construct.cpp index e1d0c04d7fdee..ffb037d439c8e 100644 --- a/clang/test/SemaOpenACC/data-construct.cpp +++ b/clang/test/SemaOpenACC/data-construct.cpp @@ -43,22 +43,22 @@ void AtLeastOneOf() { #pragma acc data default(none) ; - // expected-error@+1{{OpenACC 'data' construct must have at least one 'copy', 'copyin', 'copyout', 'create', 'no_create', 'present', 'deviceptr', 'attach' or 'default' clause}} + // expected-error@+1{{OpenACC 'data' construct must have at least one 'attach', 'copy', 'copyin', 'copyout', 'create', 'default', 'deviceptr', 'no_create', or 'present' clause}} #pragma acc data if(Var) ; - // expected-error@+1{{OpenACC 'data' construct must have at least one 'copy', 'copyin', 'copyout', 'create', 'no_create', 'present', 'deviceptr', 'attach' or 'default' clause}} + // expected-error@+1{{OpenACC 'data' construct must have at least one 'attach', 'copy', 'copyin', 'copyout', 'create', 'default', 'deviceptr', 'no_create', or 'present' clause}} #pragma acc data async ; - // expected-error@+1{{OpenACC 'data' construct must have at least one 'copy', 'copyin', 'copyout', 'create', 'no_create', 'present', 'deviceptr', 'attach' or 'default' clause}} + // expected-error@+1{{OpenACC 'data' construct must have at least one 'attach', 'copy', 'copyin', 'copyout', 'create', 'default', 'deviceptr', 'no_create', or 'present' clause}} #pragma acc data wait ; - // expected-error@+1{{OpenACC 'data' construct must have at least one 'copy', 'copyin', 'copyout', 'create', 'no_create', 'present', 'deviceptr', 'attach' or 'default' clause}} + // expected-error@+1{{OpenACC 'data' construct must have at least one 'attach', 'copy', 'copyin', 'copyout', 'create', 'default', 'deviceptr', 'no_create', or 'present' clause}} #pragma acc data device_type(*) ; - // expected-error@+1{{OpenACC 'data' construct must have at least one 'copy', 'copyin', 'copyout', 'create', 'no_create', 'present', 'deviceptr', 'attach' or 'default' clause}} + // expected-error@+1{{OpenACC 'data' construct must have at least one 'attach', 'copy', 'copyin', 'copyout', 'create', 'default', 'deviceptr', 'no_create', or 'present' clause}} #pragma acc data ; @@ -67,13 +67,13 @@ void AtLeastOneOf() { #pragma acc enter data create(Var) #pragma acc enter data attach(VarPtr) - // expected-error@+1{{OpenACC 'enter data' construct must have at least one 'copyin', 'create' or 'attach' clause}} + // expected-error@+1{{OpenACC 'enter data' construct must have at least one 'attach', 'copyin', or 'create' clause}} #pragma acc enter data if(Var) - // expected-error@+1{{OpenACC 'enter data' construct must have at least one 'copyin', 'create' or 'attach' clause}} + // expected-error@+1{{OpenACC 'enter data' construct must have at least one 'attach', 'copyin', or 'create' clause}} #pragma acc enter data async - // expected-error@+1{{OpenACC 'enter data' construct must have at least one 'copyin', 'create' or 'attach' clause}} + // expected-error@+1{{OpenACC 'enter data' construct must have at least one 'attach', 'copyin', or 'create' clause}} #pragma acc enter data wait - // expected-error@+1{{OpenACC 'enter data' construct must have at least one 'copyin', 'create' or 'attach' clause}} + // expected-error@+1{{OpenACC 'enter data' construct must have at least one 'attach', 'copyin', or 'create' clause}} #pragma acc enter data // Exit Data @@ -81,15 +81,15 @@ void AtLeastOneOf() { #pragma acc exit data delete(Var) #pragma acc exit data detach(VarPtr) - // expected-error@+1{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete' or 'detach' clause}} + // expected-error@+1{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete', or 'detach' clause}} #pragma acc exit data if(Var) - // expected-error@+1{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete' or 'detach' clause}} + // expected-error@+1{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete', or 'detach' clause}} #pragma acc exit data async - // expected-error@+1{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete' or 'detach' clause}} + // expected-error@+1{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete', or 'detach' clause}} #pragma acc exit data wait - // expected-error@+1{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete' or 'detach' clause}} + // expected-error@+1{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete', or 'detach' clause}} #pragma acc exit data finalize - // expected-error@+1{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete' or 'detach' clause}} + // expected-error@+1{{OpenACC 'exit data' construct must have at least one 'copyout', 'delete', or 'detach' clause}} #pragma acc exit data // Host Data diff --git a/clang/test/SemaOpenACC/init-construct.cpp b/clang/test/SemaOpenACC/init-construct.cpp index 2cbe5f1ded04a..ad901b3aa211e 100644 --- a/clang/test/SemaOpenACC/init-construct.cpp +++ b/clang/test/SemaOpenACC/init-construct.cpp @@ -42,7 +42,13 @@ void TestInst() { #pragma acc init #pragma acc init if (T::value < T{}) #pragma acc init device_type(radeon) device_num(getI()) if (getI() < getS()) -#pragma acc init device_type(multicore) device_type(host) device_num(t) if (t < T::value) device_num(getI()) if (getI() < getS()) + // expected-error@+2{{OpenACC 'device_num' clause cannot appear more than once on a 'init' directive}} + // expected-note@+1{{previous clause is here}} +#pragma acc init device_type(multicore) device_type(host) device_num(t) if (t < T::value) device_num(getI()) + + // expected-error@+2{{OpenACC 'if' clause cannot appear more than once on a 'init' directive}} + // expected-note@+1{{previous clause is here}} +#pragma acc init if(t < T::value) if (getI() < getS()) // expected-error@+1{{value of type 'const NotConvertible' is not contextually convertible to 'bool'}} #pragma acc init if (T::NCValue) diff --git a/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c b/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c index 2f2bd15935856..762873f13d96a 100644 --- a/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c +++ b/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c @@ -8,40 +8,40 @@ void uses() { #pragma acc loop independent for(unsigned i = 0; i < 5; ++i); - // expected-error@+2{{OpenACC clause 'seq' on 'loop' construct conflicts with previous data dependence clause}} + // expected-error@+2{{OpenACC clause 'seq' may not appear on the same construct as a 'auto' clause on a 'loop' construct}} // expected-note@+1{{previous clause is here}} #pragma acc loop auto seq for(unsigned i = 0; i < 5; ++i); - // expected-error@+2{{OpenACC clause 'independent' on 'loop' construct conflicts with previous data dependence clause}} + // expected-error@+2{{OpenACC clause 'independent' may not appear on the same construct as a 'auto' clause on a 'loop' construct}} // expected-note@+1{{previous clause is here}} #pragma acc loop auto independent for(unsigned i = 0; i < 5; ++i); - // expected-error@+2{{OpenACC clause 'auto' on 'loop' construct conflicts with previous data dependence clause}} + // expected-error@+2{{OpenACC clause 'auto' may not appear on the same construct as a 'seq' clause on a 'loop' construct}} // expected-note@+1{{previous clause is here}} #pragma acc loop seq auto for(unsigned i = 0; i < 5; ++i); - // expected-error@+2{{OpenACC clause 'independent' on 'loop' construct conflicts with previous data dependence clause}} + // expected-error@+2{{OpenACC clause 'independent' may not appear on the same construct as a 'seq' clause on a 'loop' construct}} // expected-note@+1{{previous clause is here}} #pragma acc loop seq independent for(unsigned i = 0; i < 5; ++i); - // expected-error@+2{{OpenACC clause 'auto' on 'loop' construct conflicts with previous data dependence clause}} + // expected-error@+2{{OpenACC clause 'auto' may not appear on the same construct as a 'independent' clause on a 'loop' construct}} // expected-note@+1{{previous clause is here}} #pragma acc loop independent auto for(unsigned i = 0; i < 5; ++i); - // expected-error@+2{{OpenACC clause 'seq' on 'loop' construct conflicts with previous data dependence clause}} + // expected-error@+2{{OpenACC clause 'seq' may not appear on the same construct as a 'independent' clause on a 'loop' construct}} // expected-note@+1{{previous clause is here}} #pragma acc loop independent seq for(unsigned i = 0; i < 5; ++i); - // expected-error@+2{{OpenACC clause 'seq' on 'loop' construct conflicts with previous data dependence clause}} + // expected-error@+2{{OpenACC clause 'seq' may not appear on the same construct as a 'seq' clause on a 'loop' construct}} // expected-note@+1{{previous clause is here}} #pragma acc loop seq seq for(unsigned i = 0; i < 5; ++i); - // expected-error@+2{{OpenACC clause 'independent' on 'loop' construct conflicts with previous data dependence clause}} + // expected-error@+2{{OpenACC clause 'independent' may not appear on the same construct as a 'independent' clause on a 'loop' construct}} // expected-note@+1{{previous clause is here}} #pragma acc loop independent independent for(unsigned i = 0; i < 5; ++i); - // expected-error@+2{{OpenACC clause 'auto' on 'loop' construct conflicts with previous data dependence clause}} + // expected-error@+2{{OpenACC clause 'auto' may not appear on the same construct as a 'auto' clause on a 'loop' construct}} // expected-note@+1{{previous clause is here}} #pragma acc loop auto auto for(unsigned i = 0; i < 5; ++i); @@ -591,6 +591,10 @@ void uses() { // expected-error@+2{{OpenACC clause 'gang' may not appear on the same construct as a 'seq' clause on a 'loop' construct}} // expected-note@+1{{previous clause is here}} #pragma acc loop seq gang + for(unsigned i = 0; i < 5; ++i); + // expected-error@+2{{OpenACC clause 'seq' may not appear on the same construct as a 'gang' clause on a 'loop' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc loop gang seq for(unsigned i = 0; i < 5; ++i); // expected-error@+2{{OpenACC clause 'worker' may not appear on the same construct as a 'seq' clause on a 'loop' construct}} // expected-note@+1{{previous clause is here}} diff --git a/clang/test/SemaOpenACC/loop-construct-collapse-clause.cpp b/clang/test/SemaOpenACC/loop-construct-collapse-clause.cpp index 4d45d1107d03b..851742434a1f9 100644 --- a/clang/test/SemaOpenACC/loop-construct-collapse-clause.cpp +++ b/clang/test/SemaOpenACC/loop-construct-collapse-clause.cpp @@ -18,7 +18,7 @@ void only_one_on_loop() { // expected-error@+2{{OpenACC 'collapse' clause cannot appear more than once on a 'loop' directive}} // expected-note@+1{{previous clause is here}} #pragma acc loop collapse(1) collapse(1) - for(unsigned i = 0; i < 5; ++i); + for(int i = 0; i < 5; ++i); } constexpr int three() { return 3; } @@ -323,7 +323,7 @@ void no_other_directives() { #pragma acc loop collapse(2) for(unsigned i = 0; i < 5; ++i) { for(unsigned j = 0; j < 5; ++j) { - // expected-error@+1{{OpenACC 'data' construct must have at least one 'copy', 'copyin', 'copyout', 'create', 'no_create', 'present', 'deviceptr', 'attach' or 'default' clause}} + // expected-error@+1{{OpenACC 'data' construct must have at least one 'attach', 'copy', 'copyin', 'copyout', 'create', 'default', 'deviceptr', 'no_create', or 'present' clause}} #pragma acc data ; } @@ -331,7 +331,7 @@ void no_other_directives() { // expected-note@+1{{active 'collapse' clause defined here}} #pragma acc loop collapse(2) for(unsigned i = 0; i < 5; ++i) { - // expected-error@+2{{OpenACC 'data' construct must have at least one 'copy', 'copyin', 'copyout', 'create', 'no_create', 'present', 'deviceptr', 'attach' or 'default' clause}} + // expected-error@+2{{OpenACC 'data' construct must have at least one 'attach', 'copy', 'copyin', 'copyout', 'create', 'default', 'deviceptr', 'no_create', or 'present' clause}} // expected-error@+1{{OpenACC 'data' construct cannot appear in intervening code of a 'loop' with a 'collapse' clause}} #pragma acc data for(unsigned j = 0; j < 5; ++j) { @@ -500,3 +500,43 @@ void intervening_without_force() { for(;;); } +template +void allow_multiple_collapse_templ() { + // expected-error@+2{{'collapse' clause specifies a loop count greater than the number of available loops}} + // expected-note@+1{{active 'collapse' clause defined here}} +#pragma acc loop collapse(N) device_type(*) collapse(N+1) + for(unsigned i = 0; i < 5; ++i) + for(unsigned j = 0; j < 5; ++j); + // expected-error@+2{{'collapse' clause specifies a loop count greater than the number of available loops}} + // expected-note@+1{{active 'collapse' clause defined here}} +#pragma acc loop collapse(N+1) device_type(*) collapse(N) + for(unsigned i = 0; i < 5; ++i) + for(unsigned j = 0; j < 5; ++j); +} +void allow_multiple_collapse() { + allow_multiple_collapse_templ<2>(); // expected-note{{in instantiation}} + + // expected-error@+2{{'collapse' clause specifies a loop count greater than the number of available loops}} + // expected-note@+1{{active 'collapse' clause defined here}} +#pragma acc loop collapse(2) device_type(*) collapse(3) + for(unsigned i = 0; i < 5; ++i) + for(unsigned j = 0; j < 5; ++j); + // expected-error@+2{{'collapse' clause specifies a loop count greater than the number of available loops}} + // expected-note@+1{{active 'collapse' clause defined here}} +#pragma acc loop collapse(3) device_type(*) collapse(2) + for(unsigned i = 0; i < 5; ++i) + for(unsigned j = 0; j < 5; ++j); +} + +void no_dupes_since_last_device_type() { + // expected-error@+3{{OpenACC 'collapse' clause cannot appear more than once in a 'device_type' region on a 'loop' directive}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc loop collapse(1) device_type(*) collapse(1) collapse(2) + for(unsigned i = 0; i < 5; ++i) + for(unsigned j = 0; j < 5; ++j); + +#pragma acc loop collapse(1) device_type(*) collapse(1) device_type(nvidia) collapse(2) + for(unsigned i = 0; i < 5; ++i) + for(unsigned j = 0; j < 5; ++j); +} diff --git a/clang/test/SemaOpenACC/loop-construct-tile-clause.cpp b/clang/test/SemaOpenACC/loop-construct-tile-clause.cpp index 9f6d9c3af9d9f..ebd657791f34d 100644 --- a/clang/test/SemaOpenACC/loop-construct-tile-clause.cpp +++ b/clang/test/SemaOpenACC/loop-construct-tile-clause.cpp @@ -384,6 +384,20 @@ void intervening() { } } +void use_largest_tile() { +// expected-error@+2{{'tile' clause specifies a loop count greater than the number of available loops}} +// expected-note@+1{{active 'tile' clause defined here}} +#pragma acc loop tile(1,2) device_type(*) tile (3,4,5) + for(int i = 0; i < 5; ++i) + for (int j = 0; j < 5; ++j); + +// expected-error@+2{{'tile' clause specifies a loop count greater than the number of available loops}} +// expected-note@+1{{active 'tile' clause defined here}} +#pragma acc loop tile (3,4,5) device_type(*) tile(1,2) + for(int i = 0; i < 5; ++i) + for (int j = 0; j < 5; ++j); +} + void collapse_tile_depth() { // expected-error@+4{{'collapse' clause specifies a loop count greater than the number of available loops}} // expected-note@+3{{active 'collapse' clause defined here}} @@ -394,3 +408,15 @@ void collapse_tile_depth() { for(int j = 0; j < 5; ++j); } } +void no_dupes_since_last_device_type() { + // expected-error@+3{{OpenACC 'tile' clause cannot appear more than once in a 'device_type' region on a 'loop' directive}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc loop tile(1) device_type(*) tile(1) tile(2) + for(unsigned i = 0; i < 5; ++i) + for(unsigned j = 0; j < 5; ++j); + +#pragma acc loop tile(1) device_type(*) tile(1) device_type(nvidia) tile(2) + for(unsigned i = 0; i < 5; ++i) + for(unsigned j = 0; j < 5; ++j); +} diff --git a/clang/test/SemaOpenACC/no-empty-pqr-list.cpp b/clang/test/SemaOpenACC/no-empty-pqr-list.cpp index fdac89646a5f8..eee9c96fe2c4d 100644 --- a/clang/test/SemaOpenACC/no-empty-pqr-list.cpp +++ b/clang/test/SemaOpenACC/no-empty-pqr-list.cpp @@ -213,7 +213,7 @@ void Executable(int i) { // expected-error@+1{{expected identifier}} #pragma acc shutdown device_type() // expected-error@+1{{expected identifier}} -#pragma acc set if(true) device_type() +#pragma acc set device_num(i) device_type() // expected-error@+1{{expected identifier}} #pragma acc update self(i) device_type() diff --git a/clang/test/SemaOpenACC/routine-construct-clauses.cpp b/clang/test/SemaOpenACC/routine-construct-clauses.cpp index b9fc46aeff092..004a0694c7f59 100644 --- a/clang/test/SemaOpenACC/routine-construct-clauses.cpp +++ b/clang/test/SemaOpenACC/routine-construct-clauses.cpp @@ -7,11 +7,11 @@ void Func2(); #pragma acc routine(Func) vector nohost #pragma acc routine(Func) nohost seq #pragma acc routine(Func) gang -// expected-error@+2{{OpenACC 'bind' clause cannot appear more than once on a 'routine' directive}} +// expected-error@+2{{OpenACC clause 'bind' may not appear on the same construct as a 'bind' clause on a 'routine' construct}} // expected-note@+1{{previous clause is here}} #pragma acc routine(Func) gang bind(a) bind(a) -// expected-error@+2{{OpenACC 'bind' clause cannot appear more than once on a 'routine' directive}} +// expected-error@+2{{OpenACC clause 'bind' may not appear on the same construct as a 'bind' clause on a 'routine' construct}} // expected-note@+1{{previous clause is here}} #pragma acc routine gang bind(a) bind(a) void DupeImplName(); @@ -65,9 +65,9 @@ void DupeImplName(); // expected-error@+2{{OpenACC clause 'gang' may not appear on the same construct as a 'gang' clause on a 'routine' construct}} // expected-note@+1{{previous clause is here}} #pragma acc routine(Func) gang gang -// expected-error@+1{{OpenACC 'routine' construct must have at least one 'gang', 'worker', 'vector' or 'seq' clause}} +// expected-error@+1{{OpenACC 'routine' construct must have at least one 'gang', 'seq', 'vector', or 'worker' clause}} #pragma acc routine(Func) -// expected-error@+1{{OpenACC 'routine' construct must have at least one 'gang', 'worker', 'vector' or 'seq' clause}} +// expected-error@+1{{OpenACC 'routine' construct must have at least one 'gang', 'seq', 'vector', or 'worker' clause}} #pragma acc routine(Func) nohost // only the 'dim' syntax for gang is legal. @@ -227,3 +227,548 @@ void DupeBinds6(); // expected-note@-3{{previous clause is here}} #pragma acc routine bind("asdfDiff") seq void DupeBinds6(); + +// The standard wasn't initially clear, but is being clarified that we require +// exactly one of (gang, worker, vector, or seq) to apply to each 'device_type'. +// The ones before any 'device_type' apply to all. +namespace FigureDupesAllowedAroundDeviceType { + // Test conflicts without a device type + // expected-error@+2{{OpenACC clause 'gang' may not appear on the same construct as a 'gang' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine gang gang + void Func1(); + // expected-error@+2{{OpenACC clause 'worker' may not appear on the same construct as a 'gang' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine gang worker + void Func2(); + // expected-error@+2{{OpenACC clause 'vector' may not appear on the same construct as a 'gang' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine gang vector + void Func3(); + // expected-error@+2{{OpenACC clause 'seq' may not appear on the same construct as a 'gang' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine gang seq + void Func4(); + // expected-error@+2{{OpenACC clause 'gang' may not appear on the same construct as a 'worker' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine worker gang + void Func5(); + // expected-error@+2{{OpenACC clause 'worker' may not appear on the same construct as a 'worker' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine worker worker + void Func6(); + // expected-error@+2{{OpenACC clause 'vector' may not appear on the same construct as a 'worker' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine worker vector + void Func7(); + // expected-error@+2{{OpenACC clause 'seq' may not appear on the same construct as a 'worker' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine worker seq + void Func8(); + // expected-error@+2{{OpenACC clause 'gang' may not appear on the same construct as a 'vector' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine vector gang + void Func9(); + // expected-error@+2{{OpenACC clause 'worker' may not appear on the same construct as a 'vector' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine vector worker + void Func10(); + // expected-error@+2{{OpenACC clause 'vector' may not appear on the same construct as a 'vector' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine vector vector + void Func11(); + // expected-error@+2{{OpenACC clause 'seq' may not appear on the same construct as a 'vector' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine vector seq + void Func12(); + // expected-error@+2{{OpenACC clause 'gang' may not appear on the same construct as a 'seq' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine seq gang + void Func13(); + // expected-error@+2{{OpenACC clause 'worker' may not appear on the same construct as a 'seq' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine seq worker + void Func14(); + // expected-error@+2{{OpenACC clause 'vector' may not appear on the same construct as a 'seq' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine seq vector + void Func15(); + // expected-error@+2{{OpenACC clause 'seq' may not appear on the same construct as a 'seq' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine seq seq + void Func16(); + // None included for one without a device type + // expected-error@+1{{OpenACC 'routine' construct must have at least one 'gang', 'seq', 'vector', or 'worker' clause}} +#pragma acc routine + void Func16(); + + // Check same conflicts for 'before' the device_type + // expected-error@+2{{OpenACC clause 'gang' may not appear on the same construct as a 'gang' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine gang gang device_type(*) + void Func17(); + // expected-error@+2{{OpenACC clause 'worker' may not appear on the same construct as a 'gang' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine gang worker device_type(*) + void Func18(); + // expected-error@+2{{OpenACC clause 'vector' may not appear on the same construct as a 'gang' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine gang vector device_type(*) + void Func19(); + // expected-error@+2{{OpenACC clause 'seq' may not appear on the same construct as a 'gang' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine gang seq device_type(*) + void Func20(); + // expected-error@+2{{OpenACC clause 'gang' may not appear on the same construct as a 'worker' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine worker gang device_type(*) + void Func21(); + // expected-error@+2{{OpenACC clause 'worker' may not appear on the same construct as a 'worker' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine worker worker device_type(*) + void Func22(); + // expected-error@+2{{OpenACC clause 'vector' may not appear on the same construct as a 'worker' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine worker vector device_type(*) + void Func23(); + // expected-error@+2{{OpenACC clause 'seq' may not appear on the same construct as a 'worker' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine worker seq device_type(*) + void Func24(); + // expected-error@+2{{OpenACC clause 'gang' may not appear on the same construct as a 'vector' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine vector gang device_type(*) + void Func25(); + // expected-error@+2{{OpenACC clause 'worker' may not appear on the same construct as a 'vector' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine vector worker device_type(*) + void Func26(); + // expected-error@+2{{OpenACC clause 'vector' may not appear on the same construct as a 'vector' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine vector vector device_type(*) + void Func27(); + // expected-error@+2{{OpenACC clause 'seq' may not appear on the same construct as a 'vector' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine vector seq device_type(*) + void Func28(); + // expected-error@+2{{OpenACC clause 'gang' may not appear on the same construct as a 'seq' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine seq gang device_type(*) + void Func29(); + // expected-error@+2{{OpenACC clause 'worker' may not appear on the same construct as a 'seq' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine seq worker device_type(*) + void Func30(); + // expected-error@+2{{OpenACC clause 'vector' may not appear on the same construct as a 'seq' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine seq vector device_type(*) + void Func31(); + // expected-error@+2{{OpenACC clause 'seq' may not appear on the same construct as a 'seq' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine seq seq device_type(*) + void Func32(); + // None included for the device_type + // expected-error@+1{{OpenACC 'routine' construct must have at least one 'gang', 'seq', 'vector', or 'worker' clause that applies to each 'device_type'}} +#pragma acc routine device_type(*) + void Func33(); + + // Conflicts between 'global' and 'device_type' + // expected-error@+3{{OpenACC clause 'gang' after 'device_type' clause on a 'routine' conflicts with the 'gang' clause before the first 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine gang device_type(*) gang + void Func34(); + // expected-error@+3{{OpenACC clause 'worker' after 'device_type' clause on a 'routine' conflicts with the 'gang' clause before the first 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine gang device_type(*) worker + void Func35(); + // expected-error@+3{{OpenACC clause 'vector' after 'device_type' clause on a 'routine' conflicts with the 'gang' clause before the first 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine gang device_type(*) vector + void Func36(); + // expected-error@+3{{OpenACC clause 'seq' after 'device_type' clause on a 'routine' conflicts with the 'gang' clause before the first 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine gang device_type(*) seq + void Func37(); + // expected-error@+3{{OpenACC clause 'gang' after 'device_type' clause on a 'routine' conflicts with the 'worker' clause before the first 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine worker device_type(*) gang + void Func38(); + // expected-error@+3{{OpenACC clause 'worker' after 'device_type' clause on a 'routine' conflicts with the 'worker' clause before the first 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine worker device_type(*) worker + void Func39(); + // expected-error@+3{{OpenACC clause 'vector' after 'device_type' clause on a 'routine' conflicts with the 'worker' clause before the first 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine worker device_type(*) vector + void Func40(); + // expected-error@+3{{OpenACC clause 'seq' after 'device_type' clause on a 'routine' conflicts with the 'worker' clause before the first 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine worker device_type(*) seq + void Func41(); + // expected-error@+3{{OpenACC clause 'gang' after 'device_type' clause on a 'routine' conflicts with the 'vector' clause before the first 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine vector device_type(*) gang + void Func42(); + // expected-error@+3{{OpenACC clause 'worker' after 'device_type' clause on a 'routine' conflicts with the 'vector' clause before the first 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine vector device_type(*) worker + void Func43(); + // expected-error@+3{{OpenACC clause 'vector' after 'device_type' clause on a 'routine' conflicts with the 'vector' clause before the first 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine vector device_type(*) vector + void Func44(); + // expected-error@+3{{OpenACC clause 'seq' after 'device_type' clause on a 'routine' conflicts with the 'vector' clause before the first 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine vector device_type(*) seq + void Func45(); + // expected-error@+3{{OpenACC clause 'gang' after 'device_type' clause on a 'routine' conflicts with the 'seq' clause before the first 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine seq device_type(*) gang + void Func46(); + // expected-error@+3{{OpenACC clause 'worker' after 'device_type' clause on a 'routine' conflicts with the 'seq' clause before the first 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine seq device_type(*) worker + void Func47(); + // expected-error@+3{{OpenACC clause 'vector' after 'device_type' clause on a 'routine' conflicts with the 'seq' clause before the first 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine seq device_type(*) vector + void Func48(); + // expected-error@+3{{OpenACC clause 'seq' after 'device_type' clause on a 'routine' conflicts with the 'seq' clause before the first 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine seq device_type(*) seq + void Func49(); + + // Conflicts within same device_type + // expected-error@+3{{OpenACC clause 'gang' on a 'routine' directive conflicts with the 'gang' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) gang gang + void Func50(); + // expected-error@+3{{OpenACC clause 'worker' on a 'routine' directive conflicts with the 'gang' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) gang worker + void Func51(); + // expected-error@+3{{OpenACC clause 'vector' on a 'routine' directive conflicts with the 'gang' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) gang vector + void Func52(); + // expected-error@+3{{OpenACC clause 'seq' on a 'routine' directive conflicts with the 'gang' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) gang seq + void Func53(); + // expected-error@+3{{OpenACC clause 'gang' on a 'routine' directive conflicts with the 'worker' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) worker gang + void Func54(); + // expected-error@+3{{OpenACC clause 'worker' on a 'routine' directive conflicts with the 'worker' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) worker worker + void Func55(); + // expected-error@+3{{OpenACC clause 'vector' on a 'routine' directive conflicts with the 'worker' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) worker vector + void Func56(); + // expected-error@+3{{OpenACC clause 'seq' on a 'routine' directive conflicts with the 'worker' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) worker seq + void Func57(); + // expected-error@+3{{OpenACC clause 'gang' on a 'routine' directive conflicts with the 'vector' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) vector gang + void Func58(); + // expected-error@+3{{OpenACC clause 'worker' on a 'routine' directive conflicts with the 'vector' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) vector worker + void Func59(); + // expected-error@+3{{OpenACC clause 'vector' on a 'routine' directive conflicts with the 'vector' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) vector vector + void Func60(); + // expected-error@+3{{OpenACC clause 'seq' on a 'routine' directive conflicts with the 'vector' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) vector seq + void Func61(); + // expected-error@+3{{OpenACC clause 'gang' on a 'routine' directive conflicts with the 'seq' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) seq gang + void Func62(); + // expected-error@+3{{OpenACC clause 'worker' on a 'routine' directive conflicts with the 'seq' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) seq worker + void Func63(); + // expected-error@+3{{OpenACC clause 'vector' on a 'routine' directive conflicts with the 'seq' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) seq vector + void Func64(); + // expected-error@+3{{OpenACC clause 'seq' on a 'routine' directive conflicts with the 'seq' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) seq seq + void Func65(); + + // expected-error@+3{{OpenACC clause 'gang' on a 'routine' directive conflicts with the 'gang' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) gang gang device_type(nvidia) seq + void Func66(); + // expected-error@+3{{OpenACC clause 'worker' on a 'routine' directive conflicts with the 'gang' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) gang worker device_type(nvidia) seq + void Func67(); + // expected-error@+3{{OpenACC clause 'vector' on a 'routine' directive conflicts with the 'gang' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) gang vector device_type(nvidia) seq + void Func68(); + // expected-error@+3{{OpenACC clause 'seq' on a 'routine' directive conflicts with the 'gang' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) gang seq device_type(nvidia) seq + void Func69(); + // expected-error@+3{{OpenACC clause 'gang' on a 'routine' directive conflicts with the 'worker' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) worker gang device_type(nvidia) seq + void Func70(); + // expected-error@+3{{OpenACC clause 'worker' on a 'routine' directive conflicts with the 'worker' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) worker worker device_type(nvidia) seq + void Func71(); + // expected-error@+3{{OpenACC clause 'vector' on a 'routine' directive conflicts with the 'worker' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) worker vector device_type(nvidia) seq + void Func72(); + // expected-error@+3{{OpenACC clause 'seq' on a 'routine' directive conflicts with the 'worker' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) worker seq device_type(nvidia) seq + void Func73(); + // expected-error@+3{{OpenACC clause 'gang' on a 'routine' directive conflicts with the 'vector' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) vector gang device_type(nvidia) seq + void Func74(); + // expected-error@+3{{OpenACC clause 'worker' on a 'routine' directive conflicts with the 'vector' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) vector worker device_type(nvidia) seq + void Func75(); + // expected-error@+3{{OpenACC clause 'vector' on a 'routine' directive conflicts with the 'vector' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) vector vector device_type(nvidia) seq + void Func76(); + // expected-error@+3{{OpenACC clause 'seq' on a 'routine' directive conflicts with the 'vector' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) vector seq device_type(nvidia) seq + void Func77(); + // expected-error@+3{{OpenACC clause 'gang' on a 'routine' directive conflicts with the 'seq' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) seq gang device_type(nvidia) seq + void Func78(); + // expected-error@+3{{OpenACC clause 'worker' on a 'routine' directive conflicts with the 'seq' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) seq worker device_type(nvidia) seq + void Func79(); + // expected-error@+3{{OpenACC clause 'vector' on a 'routine' directive conflicts with the 'seq' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) seq vector device_type(nvidia) seq + void Func80(); + // expected-error@+3{{OpenACC clause 'seq' on a 'routine' directive conflicts with the 'seq' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(*) seq seq device_type(nvidia) seq + void Func81(); + + // expected-error@+3{{OpenACC clause 'gang' on a 'routine' directive conflicts with the 'gang' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(nvidia) seq device_type(*) gang gang + void Func82(); + // expected-error@+3{{OpenACC clause 'worker' on a 'routine' directive conflicts with the 'gang' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(nvidia) seq device_type(*) gang worker + void Func83(); + // expected-error@+3{{OpenACC clause 'vector' on a 'routine' directive conflicts with the 'gang' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(nvidia) seq device_type(*) gang vector + void Func84(); + // expected-error@+3{{OpenACC clause 'seq' on a 'routine' directive conflicts with the 'gang' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(nvidia) seq device_type(*) gang seq + void Func85(); + // expected-error@+3{{OpenACC clause 'gang' on a 'routine' directive conflicts with the 'worker' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(nvidia) seq device_type(*) worker gang + void Func86(); + // expected-error@+3{{OpenACC clause 'worker' on a 'routine' directive conflicts with the 'worker' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(nvidia) seq device_type(*) worker worker + void Func87(); + // expected-error@+3{{OpenACC clause 'vector' on a 'routine' directive conflicts with the 'worker' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(nvidia) seq device_type(*) worker vector + void Func88(); + // expected-error@+3{{OpenACC clause 'seq' on a 'routine' directive conflicts with the 'worker' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(nvidia) seq device_type(*) worker seq + void Func89(); + // expected-error@+3{{OpenACC clause 'gang' on a 'routine' directive conflicts with the 'vector' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(nvidia) seq device_type(*) vector gang + void Func90(); + // expected-error@+3{{OpenACC clause 'worker' on a 'routine' directive conflicts with the 'vector' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(nvidia) seq device_type(*) vector worker + void Func91(); + // expected-error@+3{{OpenACC clause 'vector' on a 'routine' directive conflicts with the 'vector' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(nvidia) seq device_type(*) vector vector + void Func92(); + // expected-error@+3{{OpenACC clause 'seq' on a 'routine' directive conflicts with the 'vector' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(nvidia) seq device_type(*) vector seq + void Func93(); + // expected-error@+3{{OpenACC clause 'gang' on a 'routine' directive conflicts with the 'seq' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(nvidia) seq device_type(*) seq gang + void Func94(); + // expected-error@+3{{OpenACC clause 'worker' on a 'routine' directive conflicts with the 'seq' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(nvidia) seq device_type(*) seq worker + void Func95(); + // expected-error@+3{{OpenACC clause 'vector' on a 'routine' directive conflicts with the 'seq' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(nvidia) seq device_type(*) seq vector + void Func96(); + // expected-error@+3{{OpenACC clause 'seq' on a 'routine' directive conflicts with the 'seq' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine device_type(nvidia) seq device_type(*) seq seq + void Func97(); + + // All fine? +#pragma acc routine device_type(*) gang device_type(nvidia) gang + void Func98(); +#pragma acc routine device_type(*) gang device_type(nvidia) worker + void Func99(); +#pragma acc routine device_type(*) gang device_type(nvidia) vector + void Func100(); +#pragma acc routine device_type(*) gang device_type(nvidia) seq + void Func101(); +#pragma acc routine device_type(*) worker device_type(nvidia) gang + void Func102(); +#pragma acc routine device_type(*) worker device_type(nvidia) worker + void Func103(); +#pragma acc routine device_type(*) worker device_type(nvidia) vector + void Func104(); +#pragma acc routine device_type(*) worker device_type(nvidia) seq + void Func105(); +#pragma acc routine device_type(*) vector device_type(nvidia) gang + void Func106(); +#pragma acc routine device_type(*) vector device_type(nvidia) worker + void Func107(); +#pragma acc routine device_type(*) vector device_type(nvidia) vector + void Func108(); +#pragma acc routine device_type(*) vector device_type(nvidia) seq + void Func109(); +#pragma acc routine device_type(*) seq device_type(nvidia) gang + void Func110(); +#pragma acc routine device_type(*) seq device_type(nvidia) worker + void Func111(); +#pragma acc routine device_type(*) seq device_type(nvidia) vector + void Func112(); +#pragma acc routine device_type(*) seq device_type(nvidia) seq + void Func113(); + +} + +namespace BindDupes { + // expected-error@+2{{OpenACC clause 'bind' may not appear on the same construct as a 'bind' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine seq bind(asdf) bind(asdf) + void Func1(); + // expected-error@+2{{OpenACC clause 'bind' may not appear on the same construct as a 'bind' clause on a 'routine' construct}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine seq bind(asdf) bind(asdf) device_type(*) + void Func2(); + // expected-error@+3{{OpenACC clause 'bind' after 'device_type' clause on a 'routine' conflicts with the 'bind' clause before the first 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine seq bind(asdf) device_type(*) bind(asdf) + void Func3(); + // expected-error@+3{{OpenACC clause 'bind' after 'device_type' clause on a 'routine' conflicts with the 'bind' clause before the first 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine seq bind(asdf) device_type(*) bind(asdf) device_type(*) + void Func4(); + // expected-error@+3{{OpenACC clause 'bind' on a 'routine' directive conflicts with the 'bind' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine seq device_type(*) bind(asdf) bind(asdf) + void Func5(); + // expected-error@+3{{OpenACC clause 'bind' on a 'routine' directive conflicts with the 'bind' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine seq device_type(*) bind(asdf) bind(asdf) device_type(*) + void Func6(); + // expected-error@+3{{OpenACC clause 'bind' on a 'routine' directive conflicts with the 'bind' clause applying to the same 'device_type'}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc routine seq device_type(*) bind(asdf) device_type(*) bind(asdf) bind(asdf) + void Func7(); + +#pragma acc routine seq device_type(*) bind(asdf) device_type(*) bind(asdf) + void Func8(); +} diff --git a/clang/test/SemaOpenACC/set-construct.cpp b/clang/test/SemaOpenACC/set-construct.cpp index ac141d097532b..1a152b898ef4a 100644 --- a/clang/test/SemaOpenACC/set-construct.cpp +++ b/clang/test/SemaOpenACC/set-construct.cpp @@ -23,10 +23,10 @@ void uses() { // expected-error@+1{{value of type 'struct NotConvertible' is not contextually convertible to 'bool'}} #pragma acc set if (NC) device_type(radeon) - // expected-error@+2{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', 'device_type' or 'if' clause}} + // expected-error@+2{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', or 'device_type' clause}} // expected-error@+1{{OpenACC clause 'device_num' requires expression of integer type ('struct NotConvertible' invalid)}} #pragma acc set device_num(NC) - // expected-error@+4{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', 'device_type' or 'if' clause}} + // expected-error@+4{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', or 'device_type' clause}} // expected-error@+3{{multiple conversions from expression type 'struct AmbiguousConvert' to an integral type}} // expected-note@#AMBIG_INT{{conversion to integral type 'int'}} // expected-note@#AMBIG_SHORT{{conversion to integral type 'short'}} @@ -36,20 +36,21 @@ void uses() { #pragma acc set device_num(Explicit) // expected-error@+2{{OpenACC clause 'default_async' requires expression of integer type ('struct NotConvertible' invalid)}} - // expected-error@+1{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', 'device_type' or 'if' clause}} + // expected-error@+1{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', or 'device_type' clause}} #pragma acc set default_async(NC) // expected-error@+4{{multiple conversions from expression type 'struct AmbiguousConvert' to an integral type}} // expected-note@#AMBIG_INT{{conversion to integral type 'int'}} // expected-note@#AMBIG_SHORT{{conversion to integral type 'short'}} - // expected-error@+1{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', 'device_type' or 'if' clause}} + // expected-error@+1{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', or 'device_type' clause}} #pragma acc set default_async(Ambiguous) // expected-error@+2{{OpenACC integer expression requires explicit conversion from 'struct ExplicitConvertOnly' to 'int'}} // expected-note@#EXPL_CONV{{conversion to integral type 'int'}} #pragma acc set default_async(Explicit) - // expected-error@+1{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', 'device_type' or 'if' clause}} + // expected-error@+1{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', or 'device_type' clause}} #pragma acc set + // expected-error@+1{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', or 'device_type' clause}} #pragma acc set if (true) // expected-error@+2{{'default_async' clause cannot appear more than once on a 'set' directive}} @@ -68,10 +69,10 @@ void uses() { #pragma acc set device_type(acc_device_nvidia) if(true) if (true) // expected-error@+2{{OpenACC 'device_type' clause on a 'set' construct only permits one architecture}} - // expected-error@+1{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', 'device_type' or 'if' clause}} + // expected-error@+1{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', or 'device_type' clause}} #pragma acc set device_type(nvidia, radeon) // expected-error@+2{{OpenACC 'device_type' clause on a 'set' construct only permits one architecture}} - // expected-error@+1{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', 'device_type' or 'if' clause}} + // expected-error@+1{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', or 'device_type' clause}} #pragma acc set device_type(nonsense, nvidia, radeon) } diff --git a/clang/test/SemaOpenACC/shutdown-construct.cpp b/clang/test/SemaOpenACC/shutdown-construct.cpp index b167ac66c85d9..f13a281edcb0d 100644 --- a/clang/test/SemaOpenACC/shutdown-construct.cpp +++ b/clang/test/SemaOpenACC/shutdown-construct.cpp @@ -42,11 +42,17 @@ void TestInst() { #pragma acc shutdown #pragma acc shutdown if (T::value < T{}) #pragma acc shutdown device_type(multicore) device_num(getI()) if (getI() < getS()) -#pragma acc shutdown device_type(default) device_type(radeon) device_num(t) if (t < T::value) device_num(getI()) if (getI() < getS()) + // expected-error@+2{{OpenACC 'device_num' clause cannot appear more than once on a 'shutdown' directive}} + // expected-note@+1{{previous clause is here}} +#pragma acc shutdown device_type(default) device_type(radeon) device_num(t) if (t < T::value) device_num(getI()) // expected-error@+1{{value of type 'const NotConvertible' is not contextually convertible to 'bool'}} #pragma acc shutdown if (T::NCValue) + // expected-error@+2{{OpenACC 'if' clause cannot appear more than once on a 'shutdown' directive}} + // expected-note@+1{{previous clause is here}} +#pragma acc shutdown if(t < T::value) if (getI() < getS()) + // expected-error@+1{{OpenACC clause 'device_num' requires expression of integer type ('const NotConvertible' invalid)}} #pragma acc shutdown device_num(T::NCValue) // expected-error@+3{{multiple conversions from expression type 'const AmbiguousConvert' to an integral type}} diff --git a/clang/test/SemaOpenACC/update-construct.cpp b/clang/test/SemaOpenACC/update-construct.cpp index f0372736452df..a89dc1f20158a 100644 --- a/clang/test/SemaOpenACC/update-construct.cpp +++ b/clang/test/SemaOpenACC/update-construct.cpp @@ -36,15 +36,15 @@ void uses() { // These diagnose because there isn't at least 1 of 'self', 'host', or // 'device'. - // expected-error@+1{{OpenACC 'update' construct must have at least one 'self', 'host' or 'device' clause}} + // expected-error@+1{{OpenACC 'update' construct must have at least one 'device', 'host', or 'self' clause}} #pragma acc update async - // expected-error@+1{{OpenACC 'update' construct must have at least one 'self', 'host' or 'device' clause}} + // expected-error@+1{{OpenACC 'update' construct must have at least one 'device', 'host', or 'self' clause}} #pragma acc update wait - // expected-error@+1{{OpenACC 'update' construct must have at least one 'self', 'host' or 'device' clause}} + // expected-error@+1{{OpenACC 'update' construct must have at least one 'device', 'host', or 'self' clause}} #pragma acc update device_type(host) - // expected-error@+1{{OpenACC 'update' construct must have at least one 'self', 'host' or 'device' clause}} + // expected-error@+1{{OpenACC 'update' construct must have at least one 'device', 'host', or 'self' clause}} #pragma acc update if(true) - // expected-error@+1{{OpenACC 'update' construct must have at least one 'self', 'host' or 'device' clause}} + // expected-error@+1{{OpenACC 'update' construct must have at least one 'device', 'host', or 'self' clause}} #pragma acc update if_present // expected-error@+1{{value of type 'struct NotConvertible' is not contextually convertible to 'bool'}} diff --git a/llvm/include/llvm/Frontend/Directive/DirectiveBase.td b/llvm/include/llvm/Frontend/Directive/DirectiveBase.td index 707239c7ba364..87ce80846dd93 100644 --- a/llvm/include/llvm/Frontend/Directive/DirectiveBase.td +++ b/llvm/include/llvm/Frontend/Directive/DirectiveBase.td @@ -113,6 +113,10 @@ class Clause { // Set the prefix as optional. // `clause([prefix]: value)` bit isPrefixOptional = true; + + // When necessary because it isn't covered by rules, the name used in the + // clause in the clang::OpenACCClauseKind enum. + string clangAccSpelling = ""; } // Hold information about clause validity by version. diff --git a/llvm/include/llvm/Frontend/OpenACC/ACC.td b/llvm/include/llvm/Frontend/OpenACC/ACC.td index 898b7dc9e230f..6f6539c1ffebc 100644 --- a/llvm/include/llvm/Frontend/OpenACC/ACC.td +++ b/llvm/include/llvm/Frontend/OpenACC/ACC.td @@ -68,12 +68,14 @@ def ACCC_Copy : Clause<"copy"> { // 2.7.7 def ACCC_Copyin : Clause<"copyin"> { let flangClass = "AccObjectListWithModifier"; + let clangAccSpelling = "CopyIn"; let aliases = ["present_or_copyin", "pcopyin"]; } // 2.7.8 def ACCC_Copyout : Clause<"copyout"> { let flangClass = "AccObjectListWithModifier"; + let clangAccSpelling = "CopyOut"; let aliases = ["present_or_copyout", "pcopyout"]; } @@ -124,6 +126,7 @@ def ACCC_DeviceNum : Clause<"device_num"> { // 2.7.4 def ACCC_DevicePtr : Clause<"deviceptr"> { let flangClass = "AccObjectList"; + let clangAccSpelling = "DevicePtr"; } // 2.13.1 @@ -144,6 +147,7 @@ def ACCC_Finalize : Clause<"finalize"> {} // 2.5.14 def ACCC_FirstPrivate : Clause<"firstprivate"> { let flangClass = "AccObjectList"; + let clangAccSpelling = "FirstPrivate"; } // 2.9.2 @@ -179,7 +183,9 @@ def ACCC_NoCreate : Clause<"no_create"> { } // 2.15.1 -def ACCC_NoHost : Clause<"nohost"> {} +def ACCC_NoHost : Clause<"nohost"> { + let clangAccSpelling = "NoHost"; +} // 2.5.10 def ACCC_NumGangs : Clause<"num_gangs"> { diff --git a/llvm/include/llvm/TableGen/DirectiveEmitter.h b/llvm/include/llvm/TableGen/DirectiveEmitter.h index a2c9b2d427cce..28373271ef1e8 100644 --- a/llvm/include/llvm/TableGen/DirectiveEmitter.h +++ b/llvm/include/llvm/TableGen/DirectiveEmitter.h @@ -160,6 +160,30 @@ class Directive : public BaseRecord { } const Record *getCategory() const { return Def->getValueAsDef("category"); } + + // Clang uses a different format for names of its directives enum. + std::string getClangAccSpelling() const { + std::string Name = Def->getValueAsString("name").str(); + + // Clang calls the 'unknown' value 'invalid'. + if (Name == "unknown") + return "Invalid"; + + // Clang entries all start with a capital letter, so apply that. + Name[0] = std::toupper(Name[0]); + // Additionally, spaces/underscores are handled by capitalizing the next + // letter of the name and removing the space/underscore. + for (unsigned I = 0; I < Name.size(); ++I) { + if (Name[I] == ' ' || Name[I] == '_') { + Name.erase(I, 1); + assert(Name[I] != ' ' && Name[I] != '_' && + "No double spaces/underscores"); + Name[I] = std::toupper(Name[I]); + } + } + + return Name; + } }; // Wrapper class that contains Clause's information defined in DirectiveBase.td @@ -200,6 +224,30 @@ class Clause : public BaseRecord { return N; } + // Clang uses a different format for names of its clause enum, which can be + // overwritten with the `clangSpelling` value. So get the proper spelling + // here. + std::string getClangAccSpelling() const { + if (StringRef ClangSpelling = Def->getValueAsString("clangAccSpelling"); + !ClangSpelling.empty()) + return ClangSpelling.str(); + + std::string Name = Def->getValueAsString("name").str(); + // Clang entries all start with a capital letter, so apply that. + Name[0] = std::toupper(Name[0]); + // Additionally, underscores are handled by capitalizing the next letter of + // the name and removing the underscore. + for (unsigned I = 0; I < Name.size(); ++I) { + if (Name[I] == '_') { + Name.erase(I, 1); + assert(Name[I] != '_' && "No double underscores"); + Name[I] = std::toupper(Name[I]); + } + } + + return Name; + } + // Optional field. StringRef getEnumName() const { return Def->getValueAsString("enumClauseValue"); diff --git a/llvm/utils/TableGen/Basic/DirectiveEmitter.cpp b/llvm/utils/TableGen/Basic/DirectiveEmitter.cpp index ab68e028f1e96..b5276d695f1bc 100644 --- a/llvm/utils/TableGen/Basic/DirectiveEmitter.cpp +++ b/llvm/utils/TableGen/Basic/DirectiveEmitter.cpp @@ -759,10 +759,25 @@ static void generateGetDirectiveCategory(const DirectiveLanguage &DirLang, OS << "}\n"; } +namespace { +enum class DirectiveClauseFE { Flang, Clang }; + +StringRef getFESpelling(DirectiveClauseFE FE) { + switch (FE) { + case DirectiveClauseFE::Flang: + return "flang"; + case DirectiveClauseFE::Clang: + return "clang"; + } + llvm_unreachable("unknown FE kind"); +} +} // namespace + // Generate a simple enum set with the give clauses. static void generateClauseSet(ArrayRef Clauses, raw_ostream &OS, StringRef ClauseSetPrefix, const Directive &Dir, - const DirectiveLanguage &DirLang) { + const DirectiveLanguage &DirLang, + DirectiveClauseFE FE) { OS << "\n"; OS << " static " << DirLang.getClauseEnumSetClass() << " " << ClauseSetPrefix @@ -770,21 +785,35 @@ static void generateClauseSet(ArrayRef Clauses, raw_ostream &OS, for (const auto &C : Clauses) { VersionedClause VerClause(C); - OS << " llvm::" << DirLang.getCppNamespace() - << "::Clause::" << DirLang.getClausePrefix() - << VerClause.getClause().getFormattedName() << ",\n"; + if (FE == DirectiveClauseFE::Flang) { + OS << " llvm::" << DirLang.getCppNamespace() + << "::Clause::" << DirLang.getClausePrefix() + << VerClause.getClause().getFormattedName() << ",\n"; + } else { + assert(FE == DirectiveClauseFE::Clang); + assert(DirLang.getName() == "OpenACC"); + OS << " clang::OpenACCClauseKind::" + << VerClause.getClause().getClangAccSpelling() << ",\n"; + } } OS << " };\n"; } // Generate an enum set for the 4 kinds of clauses linked to a directive. static void generateDirectiveClauseSets(const DirectiveLanguage &DirLang, - raw_ostream &OS) { + DirectiveClauseFE FE, raw_ostream &OS) { - IfDefScope Scope("GEN_FLANG_DIRECTIVE_CLAUSE_SETS", OS); + std::string IfDefName{"GEN_"}; + IfDefName += getFESpelling(FE).upper(); + IfDefName += "_DIRECTIVE_CLAUSE_SETS"; + IfDefScope Scope(IfDefName, OS); OS << "\n"; - OS << "namespace llvm {\n"; + // The namespace has to be different for clang vs flang, as 2 structs with the + // same name but different layout is UB. So just put the 'clang' on in the + // clang namespace. + OS << "namespace " << (FE == DirectiveClauseFE::Flang ? "llvm" : "clang") + << " {\n"; // Open namespaces defined in the directive language. SmallVector Namespaces; @@ -797,13 +826,13 @@ static void generateDirectiveClauseSets(const DirectiveLanguage &DirLang, OS << " // Sets for " << Dir.getName() << "\n"; generateClauseSet(Dir.getAllowedClauses(), OS, "allowedClauses_", Dir, - DirLang); + DirLang, FE); generateClauseSet(Dir.getAllowedOnceClauses(), OS, "allowedOnceClauses_", - Dir, DirLang); + Dir, DirLang, FE); generateClauseSet(Dir.getAllowedExclusiveClauses(), OS, - "allowedExclusiveClauses_", Dir, DirLang); + "allowedExclusiveClauses_", Dir, DirLang, FE); generateClauseSet(Dir.getRequiredClauses(), OS, "requiredClauses_", Dir, - DirLang); + DirLang, FE); } // Closing namespaces @@ -817,27 +846,46 @@ static void generateDirectiveClauseSets(const DirectiveLanguage &DirLang, // The struct holds the 4 sets of enumeration for the 4 kinds of clauses // allowances (allowed, allowed once, allowed exclusive and required). static void generateDirectiveClauseMap(const DirectiveLanguage &DirLang, - raw_ostream &OS) { - - IfDefScope Scope("GEN_FLANG_DIRECTIVE_CLAUSE_MAP", OS); + DirectiveClauseFE FE, raw_ostream &OS) { + std::string IfDefName{"GEN_"}; + IfDefName += getFESpelling(FE).upper(); + IfDefName += "_DIRECTIVE_CLAUSE_MAP"; + IfDefScope Scope(IfDefName, OS); OS << "\n"; OS << "{\n"; + // The namespace has to be different for clang vs flang, as 2 structs with the + // same name but different layout is UB. So just put the 'clang' on in the + // clang namespace. + StringRef TopLevelNS = (FE == DirectiveClauseFE::Flang ? "llvm" : "clang"); + for (const Directive Dir : DirLang.getDirectives()) { - OS << " {llvm::" << DirLang.getCppNamespace() - << "::Directive::" << DirLang.getDirectivePrefix() - << Dir.getFormattedName() << ",\n"; + OS << " {"; + if (FE == DirectiveClauseFE::Flang) { + OS << TopLevelNS << "::" << DirLang.getCppNamespace() + << "::Directive::" << DirLang.getDirectivePrefix() + << Dir.getFormattedName() << ",\n"; + } else { + assert(FE == DirectiveClauseFE::Clang); + assert(DirLang.getName() == "OpenACC"); + OS << "clang::OpenACCDirectiveKind::" << Dir.getClangAccSpelling() + << ",\n"; + } + OS << " {\n"; - OS << " llvm::" << DirLang.getCppNamespace() << "::allowedClauses_" - << DirLang.getDirectivePrefix() << Dir.getFormattedName() << ",\n"; - OS << " llvm::" << DirLang.getCppNamespace() << "::allowedOnceClauses_" - << DirLang.getDirectivePrefix() << Dir.getFormattedName() << ",\n"; - OS << " llvm::" << DirLang.getCppNamespace() + OS << " " << TopLevelNS << "::" << DirLang.getCppNamespace() + << "::allowedClauses_" << DirLang.getDirectivePrefix() + << Dir.getFormattedName() << ",\n"; + OS << " " << TopLevelNS << "::" << DirLang.getCppNamespace() + << "::allowedOnceClauses_" << DirLang.getDirectivePrefix() + << Dir.getFormattedName() << ",\n"; + OS << " " << TopLevelNS << "::" << DirLang.getCppNamespace() << "::allowedExclusiveClauses_" << DirLang.getDirectivePrefix() << Dir.getFormattedName() << ",\n"; - OS << " llvm::" << DirLang.getCppNamespace() << "::requiredClauses_" - << DirLang.getDirectivePrefix() << Dir.getFormattedName() << ",\n"; + OS << " " << TopLevelNS << "::" << DirLang.getCppNamespace() + << "::requiredClauses_" << DirLang.getDirectivePrefix() + << Dir.getFormattedName() << ",\n"; OS << " }\n"; OS << " },\n"; } @@ -1072,13 +1120,24 @@ static void generateFlangClausesParser(const DirectiveLanguage &DirLang, OS << ")\n"; } +// Generate the implementation section for the enumeration in the directive +// language +static void emitDirectivesClangImpl(const DirectiveLanguage &DirLang, + raw_ostream &OS) { + // Currently we only have work to do for OpenACC, so skip otherwise. + if (DirLang.getName() != "OpenACC") + return; + + generateDirectiveClauseSets(DirLang, DirectiveClauseFE::Clang, OS); + generateDirectiveClauseMap(DirLang, DirectiveClauseFE::Clang, OS); +} // Generate the implementation section for the enumeration in the directive // language static void emitDirectivesFlangImpl(const DirectiveLanguage &DirLang, raw_ostream &OS) { - generateDirectiveClauseSets(DirLang, OS); + generateDirectiveClauseSets(DirLang, DirectiveClauseFE::Flang, OS); - generateDirectiveClauseMap(DirLang, OS); + generateDirectiveClauseMap(DirLang, DirectiveClauseFE::Flang, OS); generateFlangClauseParserClass(DirLang, OS); @@ -1211,6 +1270,8 @@ static void emitDirectivesImpl(const RecordKeeper &Records, raw_ostream &OS) { emitDirectivesFlangImpl(DirLang, OS); + emitDirectivesClangImpl(DirLang, OS); + generateClauseClassMacro(DirLang, OS); emitDirectivesBasicImpl(DirLang, OS);