Author: Erich Keane Date: 2024-10-11T09:05:19-07:00 New Revision: 5b25c31351ad1b10a3819411379b3258869c1e1b
URL: https://github.com/llvm/llvm-project/commit/5b25c31351ad1b10a3819411379b3258869c1e1b DIFF: https://github.com/llvm/llvm-project/commit/5b25c31351ad1b10a3819411379b3258869c1e1b.diff LOG: [OpenACC] Implement loop 'gang' clause. (#112006) The 'gang' clause is used to specify parallel execution of loops, thus has some complicated rules depending on the 'loop's associated compute construct. This patch implements all of those. Added: clang/test/SemaOpenACC/loop-construct-gang-ast.cpp clang/test/SemaOpenACC/loop-construct-gang-clause.cpp Modified: clang/include/clang/AST/OpenACCClause.h clang/include/clang/Basic/DiagnosticSemaKinds.td clang/include/clang/Basic/OpenACCClauses.def clang/include/clang/Basic/OpenACCKinds.h clang/include/clang/Parse/Parser.h clang/include/clang/Sema/SemaOpenACC.h clang/lib/AST/OpenACCClause.cpp clang/lib/AST/StmtProfile.cpp clang/lib/AST/TextNodeDumper.cpp clang/lib/Parse/ParseOpenACC.cpp clang/lib/Sema/SemaOpenACC.cpp clang/lib/Sema/TreeTransform.h clang/lib/Serialization/ASTReader.cpp clang/lib/Serialization/ASTWriter.cpp clang/test/AST/ast-print-openacc-loop-construct.cpp clang/test/ParserOpenACC/parse-clauses.c clang/test/SemaOpenACC/compute-construct-device_type-clause.c clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c clang/test/SemaOpenACC/loop-construct-device_type-clause.c clang/tools/libclang/CIndex.cpp Removed: ################################################################################ diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h index b500acc768e55a..f3a09eb651458d 100644 --- a/clang/include/clang/AST/OpenACCClause.h +++ b/clang/include/clang/AST/OpenACCClause.h @@ -119,32 +119,6 @@ class OpenACCSeqClause : public OpenACCClause { } }; -// Not yet implemented, but the type name is necessary for 'seq' diagnostics, so -// this provides a basic, do-nothing implementation. We still need to add this -// type to the visitors/etc, as well as get it to take its proper arguments. -class OpenACCGangClause : public OpenACCClause { -protected: - OpenACCGangClause(SourceLocation BeginLoc, SourceLocation EndLoc) - : OpenACCClause(OpenACCClauseKind::Gang, BeginLoc, EndLoc) { - llvm_unreachable("Not yet implemented"); - } - -public: - static bool classof(const OpenACCClause *C) { - return C->getClauseKind() == OpenACCClauseKind::Gang; - } - - static OpenACCGangClause * - Create(const ASTContext &Ctx, SourceLocation BeginLoc, SourceLocation EndLoc); - - child_range children() { - return child_range(child_iterator(), child_iterator()); - } - const_child_range children() const { - return const_child_range(const_child_iterator(), const_child_iterator()); - } -}; - // Not yet implemented, but the type name is necessary for 'seq' diagnostics, so // this provides a basic, do-nothing implementation. We still need to add this // type to the visitors/etc, as well as get it to take its proper arguments. @@ -177,7 +151,7 @@ class OpenACCVectorClause : public OpenACCClause { class OpenACCWorkerClause : public OpenACCClause { protected: OpenACCWorkerClause(SourceLocation BeginLoc, SourceLocation EndLoc) - : OpenACCClause(OpenACCClauseKind::Gang, BeginLoc, EndLoc) { + : OpenACCClause(OpenACCClauseKind::Worker, BeginLoc, EndLoc) { llvm_unreachable("Not yet implemented"); } @@ -535,6 +509,38 @@ class OpenACCClauseWithSingleIntExpr : public OpenACCClauseWithExprs { Expr *getIntExpr() { return hasIntExpr() ? getExprs()[0] : nullptr; }; }; +class OpenACCGangClause final + : public OpenACCClauseWithExprs, + public llvm::TrailingObjects<OpenACCGangClause, Expr *, OpenACCGangKind> { +protected: + OpenACCGangClause(SourceLocation BeginLoc, SourceLocation LParenLoc, + ArrayRef<OpenACCGangKind> GangKinds, + ArrayRef<Expr *> IntExprs, SourceLocation EndLoc); + + OpenACCGangKind getGangKind(unsigned I) const { + return getTrailingObjects<OpenACCGangKind>()[I]; + } + +public: + static bool classof(const OpenACCClause *C) { + return C->getClauseKind() == OpenACCClauseKind::Gang; + } + + size_t numTrailingObjects(OverloadToken<Expr *>) const { + return getNumExprs(); + } + + unsigned getNumExprs() const { return getExprs().size(); } + std::pair<OpenACCGangKind, const Expr *> getExpr(unsigned I) const { + return {getGangKind(I), getExprs()[I]}; + } + + static OpenACCGangClause * + Create(const ASTContext &Ctx, SourceLocation BeginLoc, + SourceLocation LParenLoc, ArrayRef<OpenACCGangKind> GangKinds, + ArrayRef<Expr *> IntExprs, SourceLocation EndLoc); +}; + class OpenACCNumWorkersClause : public OpenACCClauseWithSingleIntExpr { OpenACCNumWorkersClause(SourceLocation BeginLoc, SourceLocation LParenLoc, Expr *IntExpr, SourceLocation EndLoc); diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 41cdd09e971651..3c62a017005e59 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12576,6 +12576,7 @@ def err_acc_duplicate_clause_disallowed : Error<"OpenACC '%1' clause cannot appear more than once on a '%0' " "directive">; def note_acc_previous_clause_here : Note<"previous clause is here">; +def note_acc_previous_expr_here : Note<"previous expression is here">; def err_acc_branch_in_out_compute_construct : Error<"invalid %select{branch|return|throw}0 %select{out of|into}1 " "OpenACC Compute Construct">; @@ -12682,6 +12683,24 @@ def err_acc_insufficient_loops def err_acc_intervening_code : Error<"inner loops must be tightly nested inside a '%0' clause on " "a 'loop' construct">; +def err_acc_gang_multiple_elt + : Error<"OpenACC 'gang' clause may have at most one %select{unnamed or " + "'num'|'dim'|'static'}0 argument">; +def err_acc_gang_arg_invalid + : Error<"'%0' argument on 'gang' clause is not permitted on a%select{n " + "orphaned|||}1 'loop' construct %select{|associated with a " + "'parallel' compute construct|associated with a 'kernels' compute " + "construct|associated with a 'serial' compute construct}1">; +def err_acc_gang_dim_value + : Error<"argument to 'gang' clause dimension must be %select{a constant " + "expression|1, 2, or 3: evaluated to %1}0">; +def err_acc_gang_num_gangs_conflict + : Error<"'num' argument to 'gang' clause not allowed on a 'loop' construct " + "associated with a 'kernels' construct that has a 'num_gangs' " + "clause">; +def err_acc_gang_inside_gang + : Error<"loop with a 'gang' clause may not exist in the region of a 'gang' " + "clause on a 'kernels' compute construct">; // AMDGCN builtins diagnostics def err_amdgcn_global_load_lds_size_invalid_value : Error<"invalid size value">; diff --git a/clang/include/clang/Basic/OpenACCClauses.def b/clang/include/clang/Basic/OpenACCClauses.def index a380e5ae69c418..2a098de31eb618 100644 --- a/clang/include/clang/Basic/OpenACCClauses.def +++ b/clang/include/clang/Basic/OpenACCClauses.def @@ -42,6 +42,7 @@ VISIT_CLAUSE(DevicePtr) VISIT_CLAUSE(DeviceType) CLAUSE_ALIAS(DType, DeviceType, false) VISIT_CLAUSE(FirstPrivate) +VISIT_CLAUSE(Gang) VISIT_CLAUSE(If) VISIT_CLAUSE(Independent) VISIT_CLAUSE(NoCreate) diff --git a/clang/include/clang/Basic/OpenACCKinds.h b/clang/include/clang/Basic/OpenACCKinds.h index c4dfe3bedc13a7..3f48ebca708a42 100644 --- a/clang/include/clang/Basic/OpenACCKinds.h +++ b/clang/include/clang/Basic/OpenACCKinds.h @@ -550,6 +550,35 @@ inline llvm::raw_ostream &operator<<(llvm::raw_ostream &Out, OpenACCReductionOperator Op) { return printOpenACCReductionOperator(Out, Op); } + +enum class OpenACCGangKind : uint8_t { + /// num: + Num, + /// dim: + Dim, + /// static: + Static +}; + +template <typename StreamTy> +inline StreamTy &printOpenACCGangKind(StreamTy &Out, OpenACCGangKind GK) { + switch (GK) { + case OpenACCGangKind::Num: + return Out << "num"; + case OpenACCGangKind::Dim: + return Out << "dim"; + case OpenACCGangKind::Static: + return Out << "static"; + } +} +inline const StreamingDiagnostic &operator<<(const StreamingDiagnostic &Out, + OpenACCGangKind Op) { + return printOpenACCGangKind(Out, Op); +} +inline llvm::raw_ostream &operator<<(llvm::raw_ostream &Out, + OpenACCGangKind Op) { + return printOpenACCGangKind(Out, Op); +} } // namespace clang #endif // LLVM_CLANG_BASIC_OPENACCKINDS_H diff --git a/clang/include/clang/Parse/Parser.h b/clang/include/clang/Parse/Parser.h index dbcb545058a026..045ee754a242b3 100644 --- a/clang/include/clang/Parse/Parser.h +++ b/clang/include/clang/Parse/Parser.h @@ -3797,9 +3797,15 @@ class Parser : public CodeCompletionHandler { bool ParseOpenACCSizeExprList(OpenACCClauseKind CK, llvm::SmallVectorImpl<Expr *> &SizeExprs); /// Parses a 'gang-arg-list', used for the 'gang' clause. - bool ParseOpenACCGangArgList(SourceLocation GangLoc); - /// Parses a 'gang-arg', used for the 'gang' clause. - bool ParseOpenACCGangArg(SourceLocation GangLoc); + bool ParseOpenACCGangArgList(SourceLocation GangLoc, + llvm::SmallVectorImpl<OpenACCGangKind> &GKs, + llvm::SmallVectorImpl<Expr *> &IntExprs); + + using OpenACCGangArgRes = std::pair<OpenACCGangKind, ExprResult>; + /// Parses a 'gang-arg', used for the 'gang' clause. Returns a pair of the + /// ExprResult (which contains the validity of the expression), plus the gang + /// kind for the current argument. + OpenACCGangArgRes ParseOpenACCGangArg(SourceLocation GangLoc); /// Parses a 'condition' expr, ensuring it results in a ExprResult ParseOpenACCConditionExpr(); diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h index 97386d2378b758..59a9648d5f9380 100644 --- a/clang/include/clang/Sema/SemaOpenACC.h +++ b/clang/include/clang/Sema/SemaOpenACC.h @@ -38,9 +38,20 @@ class SemaOpenACC : public SemaBase { /// haven't had their 'parent' compute construct set yet. Entires will only be /// made to this list in the case where we know the loop isn't an orphan. llvm::SmallVector<OpenACCLoopConstruct *> ParentlessLoopConstructs; - /// Whether we are inside of a compute construct, and should add loops to the - /// above collection. - bool InsideComputeConstruct = false; + + struct ComputeConstructInfo { + /// Which type of compute construct we are inside of, which we can use to + /// determine whether we should add loops to the above collection. We can + /// also use it to diagnose loop construct clauses. + OpenACCDirectiveKind Kind = OpenACCDirectiveKind::Invalid; + // If we have an active compute construct, stores the list of clauses we've + // prepared for it, so that we can diagnose limitations on child constructs. + ArrayRef<OpenACCClause *> Clauses; + } ActiveComputeConstructInfo; + + bool isInComputeConstruct() const { + return ActiveComputeConstructInfo.Kind != OpenACCDirectiveKind::Invalid; + } /// Certain clauses care about the same things that aren't specific to the /// individual clause, but can be shared by a few, so store them here. All @@ -99,6 +110,15 @@ class SemaOpenACC : public SemaBase { } TileInfo; public: + ComputeConstructInfo &getActiveComputeConstructInfo() { + return ActiveComputeConstructInfo; + } + + /// If there is a current 'active' loop construct with a 'gang' clause on a + /// 'kernel' construct, this will have the source location for it. This + /// permits us to implement the restriction of no further 'gang' clauses. + SourceLocation LoopGangClauseOnKernelLoc; + // Redeclaration of the version in OpenACCClause.h. using DeviceTypeArgument = std::pair<IdentifierInfo *, SourceLocation>; @@ -149,9 +169,14 @@ class SemaOpenACC : public SemaBase { Expr *LoopCount; }; + struct GangDetails { + SmallVector<OpenACCGangKind> GangKinds; + SmallVector<Expr *> IntExprs; + }; + std::variant<std::monostate, DefaultDetails, ConditionDetails, IntExprDetails, VarListDetails, WaitDetails, DeviceTypeDetails, - ReductionDetails, CollapseDetails> + ReductionDetails, CollapseDetails, GangDetails> Details = std::monostate{}; public: @@ -245,9 +270,18 @@ class SemaOpenACC : public SemaBase { ClauseKind == OpenACCClauseKind::NumWorkers || ClauseKind == OpenACCClauseKind::Async || ClauseKind == OpenACCClauseKind::Tile || + ClauseKind == OpenACCClauseKind::Gang || ClauseKind == OpenACCClauseKind::VectorLength) && "Parsed clause kind does not have a int exprs"); + if (ClauseKind == OpenACCClauseKind::Gang) { + // There might not be any gang int exprs, as this is an optional + // argument. + if (std::holds_alternative<std::monostate>(Details)) + return {}; + return std::get<GangDetails>(Details).IntExprs; + } + return std::get<IntExprDetails>(Details).IntExprs; } @@ -259,6 +293,16 @@ class SemaOpenACC : public SemaBase { return std::get<ReductionDetails>(Details).Op; } + ArrayRef<OpenACCGangKind> getGangKinds() const { + assert(ClauseKind == OpenACCClauseKind::Gang && + "Parsed clause kind does not have gang kind"); + // The args on gang are optional, so this might not actually hold + // anything. + if (std::holds_alternative<std::monostate>(Details)) + return {}; + return std::get<GangDetails>(Details).GangKinds; + } + ArrayRef<Expr *> getVarList() { assert((ClauseKind == OpenACCClauseKind::Private || ClauseKind == OpenACCClauseKind::NoCreate || @@ -371,6 +415,25 @@ class SemaOpenACC : public SemaBase { Details = IntExprDetails{std::move(IntExprs)}; } + void setGangDetails(ArrayRef<OpenACCGangKind> GKs, + ArrayRef<Expr *> IntExprs) { + assert(ClauseKind == OpenACCClauseKind::Gang && + "Parsed Clause kind does not have gang details"); + assert(GKs.size() == IntExprs.size() && "Mismatched kind/size?"); + + Details = GangDetails{{GKs.begin(), GKs.end()}, + {IntExprs.begin(), IntExprs.end()}}; + } + + void setGangDetails(llvm::SmallVector<OpenACCGangKind> &&GKs, + llvm::SmallVector<Expr *> &&IntExprs) { + assert(ClauseKind == OpenACCClauseKind::Gang && + "Parsed Clause kind does not have gang details"); + assert(GKs.size() == IntExprs.size() && "Mismatched kind/size?"); + + Details = GangDetails{std::move(GKs), std::move(IntExprs)}; + } + void setVarListDetails(ArrayRef<Expr *> VarList, bool IsReadOnly, bool IsZero) { assert((ClauseKind == OpenACCClauseKind::Private || @@ -545,10 +608,12 @@ class SemaOpenACC : public SemaBase { SourceLocation RBLoc); /// Checks the loop depth value for a collapse clause. ExprResult CheckCollapseLoopCount(Expr *LoopCount); - /// Checks a single size expr for a tile clause. 'gang' could possibly call - /// this, but has slightly stricter rules as to valid values. + /// Checks a single size expr for a tile clause. ExprResult CheckTileSizeExpr(Expr *SizeExpr); + // Check a single expression on a gang clause. + ExprResult CheckGangExpr(OpenACCGangKind GK, Expr *E); + ExprResult BuildOpenACCAsteriskSizeExpr(SourceLocation AsteriskLoc); ExprResult ActOnOpenACCAsteriskSizeExpr(SourceLocation AsteriskLoc); @@ -595,8 +660,9 @@ class SemaOpenACC : public SemaBase { /// Loop needing its parent construct. class AssociatedStmtRAII { SemaOpenACC &SemaRef; - bool WasInsideComputeConstruct; + ComputeConstructInfo OldActiveComputeConstructInfo; OpenACCDirectiveKind DirKind; + SourceLocation OldLoopGangClauseOnKernelLoc; llvm::SmallVector<OpenACCLoopConstruct *> ParentlessLoopConstructs; LoopInConstructRAII LoopRAII; diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp index 0b34ed6189593e..6fb8fe0b8cfeef 100644 --- a/clang/lib/AST/OpenACCClause.cpp +++ b/clang/lib/AST/OpenACCClause.cpp @@ -26,7 +26,7 @@ bool OpenACCClauseWithExprs::classof(const OpenACCClause *C) { return OpenACCWaitClause::classof(C) || OpenACCNumGangsClause::classof(C) || OpenACCTileClause::classof(C) || OpenACCClauseWithSingleIntExpr::classof(C) || - OpenACCClauseWithVarList::classof(C); + OpenACCGangClause::classof(C) || OpenACCClauseWithVarList::classof(C); } bool OpenACCClauseWithVarList::classof(const OpenACCClause *C) { return OpenACCPrivateClause::classof(C) || @@ -125,6 +125,21 @@ OpenACCNumWorkersClause::OpenACCNumWorkersClause(SourceLocation BeginLoc, "Condition expression type not scalar/dependent"); } +OpenACCGangClause::OpenACCGangClause(SourceLocation BeginLoc, + SourceLocation LParenLoc, + ArrayRef<OpenACCGangKind> GangKinds, + ArrayRef<Expr *> IntExprs, + SourceLocation EndLoc) + : OpenACCClauseWithExprs(OpenACCClauseKind::Gang, BeginLoc, LParenLoc, + EndLoc) { + assert(GangKinds.size() == IntExprs.size() && "Mismatch exprs/kind?"); + std::uninitialized_copy(IntExprs.begin(), IntExprs.end(), + getTrailingObjects<Expr *>()); + setExprs(MutableArrayRef(getTrailingObjects<Expr *>(), IntExprs.size())); + std::uninitialized_copy(GangKinds.begin(), GangKinds.end(), + getTrailingObjects<OpenACCGangKind>()); +} + OpenACCNumWorkersClause * OpenACCNumWorkersClause::Create(const ASTContext &C, SourceLocation BeginLoc, SourceLocation LParenLoc, Expr *IntExpr, @@ -376,11 +391,16 @@ OpenACCSeqClause *OpenACCSeqClause::Create(const ASTContext &C, return new (Mem) OpenACCSeqClause(BeginLoc, EndLoc); } -OpenACCGangClause *OpenACCGangClause::Create(const ASTContext &C, - SourceLocation BeginLoc, - SourceLocation EndLoc) { - void *Mem = C.Allocate(sizeof(OpenACCGangClause)); - return new (Mem) OpenACCGangClause(BeginLoc, EndLoc); +OpenACCGangClause * +OpenACCGangClause::Create(const ASTContext &C, SourceLocation BeginLoc, + SourceLocation LParenLoc, + ArrayRef<OpenACCGangKind> GangKinds, + ArrayRef<Expr *> IntExprs, SourceLocation EndLoc) { + void *Mem = + C.Allocate(OpenACCGangClause::totalSizeToAlloc<Expr *, OpenACCGangKind>( + IntExprs.size(), GangKinds.size())); + return new (Mem) + OpenACCGangClause(BeginLoc, LParenLoc, GangKinds, IntExprs, EndLoc); } OpenACCWorkerClause *OpenACCWorkerClause::Create(const ASTContext &C, @@ -600,3 +620,21 @@ void OpenACCClausePrinter::VisitCollapseClause(const OpenACCCollapseClause &C) { printExpr(C.getLoopCount()); OS << ")"; } + +void OpenACCClausePrinter::VisitGangClause(const OpenACCGangClause &C) { + OS << "gang"; + + if (C.getNumExprs() > 0) { + OS << "("; + bool first = true; + for (unsigned I = 0; I < C.getNumExprs(); ++I) { + if (!first) + OS << ", "; + first = false; + + OS << C.getExpr(I).first << ": "; + printExpr(C.getExpr(I).second); + } + OS << ")"; + } +} diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp index 4d177fd6c5968c..6161b1403ed35d 100644 --- a/clang/lib/AST/StmtProfile.cpp +++ b/clang/lib/AST/StmtProfile.cpp @@ -2646,6 +2646,12 @@ void OpenACCClauseProfiler::VisitIndependentClause( void OpenACCClauseProfiler::VisitSeqClause(const OpenACCSeqClause &Clause) {} +void OpenACCClauseProfiler::VisitGangClause(const OpenACCGangClause &Clause) { + for (unsigned I = 0; I < Clause.getNumExprs(); ++I) { + Profiler.VisitStmt(Clause.getExpr(I).second); + } +} + void OpenACCClauseProfiler::VisitReductionClause( const OpenACCReductionClause &Clause) { for (auto *E : Clause.getVarList()) diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp index 15b23d60c3ffab..ac8c196777f9b8 100644 --- a/clang/lib/AST/TextNodeDumper.cpp +++ b/clang/lib/AST/TextNodeDumper.cpp @@ -425,6 +425,17 @@ void TextNodeDumper::Visit(const OpenACCClause *C) { // but print 'clause' here so it is clear what is happening from the dump. OS << " clause"; break; + case OpenACCClauseKind::Gang: { + OS << " clause"; + // print the list of all GangKinds, so that there is some sort of + // relationship to the expressions listed afterwards. + auto *GC = cast<OpenACCGangClause>(C); + + for (unsigned I = 0; I < GC->getNumExprs(); ++I) { + OS << " " << GC->getExpr(I).first; + } + break; + } case OpenACCClauseKind::Collapse: OS << " clause"; if (cast<OpenACCCollapseClause>(C)->hasForce()) diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp index b27e50b147f4a8..635039b724e6a0 100644 --- a/clang/lib/Parse/ParseOpenACC.cpp +++ b/clang/lib/Parse/ParseOpenACC.cpp @@ -797,23 +797,26 @@ bool Parser::ParseOpenACCSizeExprList( /// [num:]int-expr /// dim:int-expr /// static:size-expr -bool Parser::ParseOpenACCGangArg(SourceLocation GangLoc) { +Parser::OpenACCGangArgRes Parser::ParseOpenACCGangArg(SourceLocation GangLoc) { if (isOpenACCSpecialToken(OpenACCSpecialTokenKind::Static, getCurToken()) && NextToken().is(tok::colon)) { // 'static' just takes a size-expr, which is an int-expr or an asterisk. ConsumeToken(); ConsumeToken(); - return ParseOpenACCSizeExpr(OpenACCClauseKind::Gang).isInvalid(); + ExprResult Res = ParseOpenACCSizeExpr(OpenACCClauseKind::Gang); + return {OpenACCGangKind::Static, Res}; } if (isOpenACCSpecialToken(OpenACCSpecialTokenKind::Dim, getCurToken()) && NextToken().is(tok::colon)) { ConsumeToken(); ConsumeToken(); - return ParseOpenACCIntExpr(OpenACCDirectiveKind::Invalid, - OpenACCClauseKind::Gang, GangLoc) - .first.isInvalid(); + // Parse this as a const-expression, and we'll check its integer-ness/value + // in CheckGangExpr. + ExprResult Res = + getActions().CorrectDelayedTyposInExpr(ParseConstantExpression()); + return {OpenACCGangKind::Dim, Res}; } if (isOpenACCSpecialToken(OpenACCSpecialTokenKind::Num, getCurToken()) && @@ -822,27 +825,40 @@ bool Parser::ParseOpenACCGangArg(SourceLocation GangLoc) { ConsumeToken(); // Fallthrough to the 'int-expr' handling for when 'num' is omitted. } + // This is just the 'num' case where 'num' is optional. - return ParseOpenACCIntExpr(OpenACCDirectiveKind::Invalid, - OpenACCClauseKind::Gang, GangLoc) - .first.isInvalid(); + ExprResult Res = ParseOpenACCIntExpr(OpenACCDirectiveKind::Invalid, + OpenACCClauseKind::Gang, GangLoc) + .first; + return {OpenACCGangKind::Num, Res}; } -bool Parser::ParseOpenACCGangArgList(SourceLocation GangLoc) { - if (ParseOpenACCGangArg(GangLoc)) { +bool Parser::ParseOpenACCGangArgList( + SourceLocation GangLoc, llvm::SmallVectorImpl<OpenACCGangKind> &GKs, + llvm::SmallVectorImpl<Expr *> &IntExprs) { + + Parser::OpenACCGangArgRes Res = ParseOpenACCGangArg(GangLoc); + if (!Res.second.isUsable()) { SkipUntil(tok::r_paren, tok::annot_pragma_openacc_end, Parser::StopBeforeMatch); - return false; + return true; } + GKs.push_back(Res.first); + IntExprs.push_back(Res.second.get()); + while (!getCurToken().isOneOf(tok::r_paren, tok::annot_pragma_openacc_end)) { ExpectAndConsume(tok::comma); - if (ParseOpenACCGangArg(GangLoc)) { + Res = ParseOpenACCGangArg(GangLoc); + if (!Res.second.isUsable()) { SkipUntil(tok::r_paren, tok::annot_pragma_openacc_end, Parser::StopBeforeMatch); - return false; + return true; } + + GKs.push_back(Res.first); + IntExprs.push_back(Res.second.get()); } return false; } @@ -1129,12 +1145,16 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams( } break; } - case OpenACCClauseKind::Gang: - if (ParseOpenACCGangArgList(ClauseLoc)) { + case OpenACCClauseKind::Gang: { + llvm::SmallVector<OpenACCGangKind> GKs; + llvm::SmallVector<Expr *> IntExprs; + if (ParseOpenACCGangArgList(ClauseLoc, GKs, IntExprs)) { Parens.skipToEnd(); return OpenACCCanContinue(); } + ParsedClause.setGangDetails(std::move(GKs), std::move(IntExprs)); break; + } case OpenACCClauseKind::Wait: { OpenACCWaitParseInfo Info = ParseOpenACCWaitArgument(ClauseLoc, diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp index 66f8029a2754b9..30d73d621db69b 100644 --- a/clang/lib/Sema/SemaOpenACC.cpp +++ b/clang/lib/Sema/SemaOpenACC.cpp @@ -366,6 +366,19 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind, } } + 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; + } + } + default: // Do nothing so we can go to the 'unimplemented' diagnostic instead. return true; @@ -459,6 +472,23 @@ class SemaOpenACCClauseVisitor { return nullptr; } + // OpenACC 3.3 2.9: + // A 'gang', 'worker', or 'vector' clause may not appear if a 'seq' clause + // appears. + bool DiagIfSeqClause(SemaOpenACC::OpenACCParsedClause &Clause) { + const auto *Itr = + llvm::find_if(ExistingClauses, llvm::IsaPred<OpenACCSeqClause>); + + if (Itr != ExistingClauses.end()) { + SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_cannot_combine) + << Clause.getClauseKind() << (*Itr)->getClauseKind(); + SemaRef.Diag((*Itr)->getBeginLoc(), diag::note_acc_previous_clause_here); + + return true; + } + return false; + } + public: SemaOpenACCClauseVisitor(SemaOpenACC &S, ArrayRef<const OpenACCClause *> ExistingClauses) @@ -470,26 +500,14 @@ class SemaOpenACCClauseVisitor { OpenACCClause *Visit(SemaOpenACC::OpenACCParsedClause &Clause) { switch (Clause.getClauseKind()) { - case OpenACCClauseKind::Gang: - case OpenACCClauseKind::Worker: - case OpenACCClauseKind::Vector: { - // TODO OpenACC: These are only implemented enough for the 'seq' diagnostic, - // otherwise treats itself as unimplemented. When we implement these, we - // can remove them from here. - - // OpenACC 3.3 2.9: - // A 'gang', 'worker', or 'vector' clause may not appear if a 'seq' clause - // appears. - const auto *Itr = - llvm::find_if(ExistingClauses, llvm::IsaPred<OpenACCSeqClause>); - - if (Itr != ExistingClauses.end()) { - SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_cannot_combine) - << Clause.getClauseKind() << (*Itr)->getClauseKind(); - SemaRef.Diag((*Itr)->getBeginLoc(), diag::note_acc_previous_clause_here); + case OpenACCClauseKind::Worker: + case OpenACCClauseKind::Vector: { + // TODO OpenACC: These are only implemented enough for the 'seq' + // diagnostic, otherwise treats itself as unimplemented. When we + // implement these, we can remove them from here. + DiagIfSeqClause(Clause); + return isNotImplemented(); } - return isNotImplemented(); - } #define VISIT_CLAUSE(CLAUSE_NAME) \ case OpenACCClauseKind::CLAUSE_NAME: \ @@ -1006,6 +1024,84 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitIndependentClause( Clause.getEndLoc()); } +OpenACCClause *SemaOpenACCClauseVisitor::VisitGangClause( + SemaOpenACC::OpenACCParsedClause &Clause) { + if (DiagIfSeqClause(Clause)) + return nullptr; + + // Restrictions only properly implemented on 'loop' constructs, and it is + // the only construct that can do anything with this, so skip/treat as + // unimplemented for the combined constructs. + if (Clause.getDirectiveKind() != OpenACCDirectiveKind::Loop) + return isNotImplemented(); + + llvm::SmallVector<OpenACCGangKind> GangKinds; + llvm::SmallVector<Expr *> IntExprs; + + // Store the existing locations, so we can do duplicate checking. Index is + // the int-value of the OpenACCGangKind enum. + SourceLocation ExistingElemLoc[3]; + + for (unsigned I = 0; I < Clause.getIntExprs().size(); ++I) { + OpenACCGangKind GK = Clause.getGangKinds()[I]; + ExprResult ER = SemaRef.CheckGangExpr(GK, Clause.getIntExprs()[I]); + + if (!ER.isUsable()) + continue; + + // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels + // construct, the gang clause behaves as follows. ... An argument with no + // keyword or with num keyword is only allowed when num_gangs does not + // appear on the kernels construct. + if (SemaRef.getActiveComputeConstructInfo().Kind == + OpenACCDirectiveKind::Kernels && + GK == OpenACCGangKind::Num) { + const auto *Itr = + llvm::find_if(SemaRef.getActiveComputeConstructInfo().Clauses, + llvm::IsaPred<OpenACCNumGangsClause>); + + if (Itr != SemaRef.getActiveComputeConstructInfo().Clauses.end()) { + SemaRef.Diag(ER.get()->getBeginLoc(), + diag::err_acc_gang_num_gangs_conflict); + SemaRef.Diag((*Itr)->getBeginLoc(), + diag::note_acc_previous_clause_here); + continue; + } + } + + // OpenACC 3.3 2.9: 'gang-arg-list' may have at most one num, one dim, and + // one static argument. + if (ExistingElemLoc[static_cast<unsigned>(GK)].isValid()) { + SemaRef.Diag(ER.get()->getBeginLoc(), diag::err_acc_gang_multiple_elt) + << static_cast<unsigned>(GK); + SemaRef.Diag(ExistingElemLoc[static_cast<unsigned>(GK)], + diag::note_acc_previous_expr_here); + continue; + } + + ExistingElemLoc[static_cast<unsigned>(GK)] = ER.get()->getBeginLoc(); + GangKinds.push_back(GK); + IntExprs.push_back(ER.get()); + } + + // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels + // construct, the gang clause behaves as follows. ... The region of a loop + // with a gang clause may not contain another loop with a gang clause unless + // within a nested compute region. + if (SemaRef.LoopGangClauseOnKernelLoc.isValid()) { + // This handles the 'inner loop' diagnostic, but we cannot set that we're on + // one of these until we get to the end of the construct. + SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_gang_inside_gang); + SemaRef.Diag(SemaRef.LoopGangClauseOnKernelLoc, + diag::note_acc_previous_clause_here); + return nullptr; + } + + return OpenACCGangClause::Create(Ctx, Clause.getBeginLoc(), + Clause.getLParenLoc(), GangKinds, IntExprs, + Clause.getEndLoc()); +} + OpenACCClause *SemaOpenACCClauseVisitor::VisitSeqClause( SemaOpenACC::OpenACCParsedClause &Clause) { // Restrictions only properly implemented on 'loop' constructs, and it is @@ -1118,17 +1214,44 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII( SemaOpenACC &S, OpenACCDirectiveKind DK, ArrayRef<const OpenACCClause *> UnInstClauses, ArrayRef<OpenACCClause *> Clauses) - : SemaRef(S), WasInsideComputeConstruct(S.InsideComputeConstruct), - DirKind(DK), LoopRAII(SemaRef, /*PreserveDepth=*/false) { + : SemaRef(S), OldActiveComputeConstructInfo(S.ActiveComputeConstructInfo), + DirKind(DK), OldLoopGangClauseOnKernelLoc(S.LoopGangClauseOnKernelLoc), + LoopRAII(SemaRef, /*PreserveDepth=*/false) { // Compute constructs end up taking their 'loop'. if (DirKind == OpenACCDirectiveKind::Parallel || DirKind == OpenACCDirectiveKind::Serial || DirKind == OpenACCDirectiveKind::Kernels) { - SemaRef.InsideComputeConstruct = true; + SemaRef.ActiveComputeConstructInfo.Kind = DirKind; + SemaRef.ActiveComputeConstructInfo.Clauses = Clauses; SemaRef.ParentlessLoopConstructs.swap(ParentlessLoopConstructs); + + // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels + // construct, the gang clause behaves as follows. ... The region of a loop + // with a gang clause may not contain another loop with a gang clause unless + // within a nested compute region. + // + // Implement the 'unless within a nested compute region' part. + SemaRef.LoopGangClauseOnKernelLoc = {}; } else if (DirKind == OpenACCDirectiveKind::Loop) { SetCollapseInfoBeforeAssociatedStmt(UnInstClauses, Clauses); SetTileInfoBeforeAssociatedStmt(UnInstClauses, Clauses); + + // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels + // construct, the gang clause behaves as follows. ... The region of a loop + // with a gang clause may not contain another loop with a gang clause unless + // within a nested compute region. + // + // We don't bother doing this when this is a template instantiation, as + // there is no reason to do these checks: the existance of a + // gang/kernels/etc cannot be dependent. + if (SemaRef.getActiveComputeConstructInfo().Kind == + OpenACCDirectiveKind::Kernels && + UnInstClauses.empty()) { + // This handles the 'outer loop' part of this. + auto *Itr = llvm::find_if(Clauses, llvm::IsaPred<OpenACCGangClause>); + if (Itr != Clauses.end()) + SemaRef.LoopGangClauseOnKernelLoc = (*Itr)->getBeginLoc(); + } } } @@ -1199,7 +1322,9 @@ void SemaOpenACC::AssociatedStmtRAII::SetTileInfoBeforeAssociatedStmt( } SemaOpenACC::AssociatedStmtRAII::~AssociatedStmtRAII() { - SemaRef.InsideComputeConstruct = WasInsideComputeConstruct; + SemaRef.ActiveComputeConstructInfo = OldActiveComputeConstructInfo; + SemaRef.LoopGangClauseOnKernelLoc = OldLoopGangClauseOnKernelLoc; + if (DirKind == OpenACCDirectiveKind::Parallel || DirKind == OpenACCDirectiveKind::Serial || DirKind == OpenACCDirectiveKind::Kernels) { @@ -1761,6 +1886,109 @@ ExprResult SemaOpenACC::CheckCollapseLoopCount(Expr *LoopCount) { ConstantExpr::Create(getASTContext(), LoopCount, APValue{*ICE})}; } +namespace { +ExprResult CheckGangStaticExpr(SemaOpenACC &S, Expr *E) { + if (isa<OpenACCAsteriskSizeExpr>(E)) + return E; + return S.ActOnIntExpr(OpenACCDirectiveKind::Invalid, OpenACCClauseKind::Gang, + E->getBeginLoc(), E); +} +} // namespace + +ExprResult SemaOpenACC::CheckGangExpr(OpenACCGangKind GK, Expr *E) { + // Gang Expr legality depends on the associated compute construct. + switch (ActiveComputeConstructInfo.Kind) { + case OpenACCDirectiveKind::Invalid: + case OpenACCDirectiveKind::Parallel: { + switch (GK) { + // OpenACC 3.3 2.9.2: When the parent compute construct is a parallel + // construct, or an orphaned loop construct, the gang clause behaves as + // follows. ... The dim argument must be a constant positive integer value + // 1, 2, or 3. + case OpenACCGangKind::Dim: { + if (!E) + return ExprError(); + ExprResult Res = + ActOnIntExpr(OpenACCDirectiveKind::Invalid, OpenACCClauseKind::Gang, + E->getBeginLoc(), E); + + if (!Res.isUsable()) + return Res; + + if (Res.get()->isInstantiationDependent()) + return Res; + + std::optional<llvm::APSInt> ICE = + Res.get()->getIntegerConstantExpr(getASTContext()); + + if (!ICE || *ICE <= 0 || ICE > 3) { + Diag(Res.get()->getBeginLoc(), diag::err_acc_gang_dim_value) + << ICE.has_value() << ICE.value_or(llvm::APSInt{}).getExtValue(); + return ExprError(); + } + + return ExprResult{ + ConstantExpr::Create(getASTContext(), Res.get(), APValue{*ICE})}; + } + // OpenACC 3.3 2.9.2: When the parent compute construct is a parallel + // construct, or an orphaned loop construct, the gang clause behaves as + // follows. ... The num argument is not allowed. + case OpenACCGangKind::Num: + Diag(E->getBeginLoc(), diag::err_acc_gang_arg_invalid) + << GK + << (/*orphan/parallel=*/ActiveComputeConstructInfo.Kind == + OpenACCDirectiveKind::Parallel + ? 1 + : 0); + return ExprError(); + case OpenACCGangKind::Static: + return CheckGangStaticExpr(*this, E); + } + } break; + case OpenACCDirectiveKind::Kernels: { + switch (GK) { + // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels + // construct, the gang clause behaves as follows. ... The dim argument is + // not allowed. + case OpenACCGangKind::Dim: + Diag(E->getBeginLoc(), diag::err_acc_gang_arg_invalid) + << GK << /*kernels=*/2; + return ExprError(); + // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels + // construct, the gang clause behaves as follows. ... An argument with no + // keyword or with num keyword is only allowed when num_gangs does not + // appear on the kernels construct. ... The region of a loop with the gang + // clause may not contain another loop with a gang clause unless within a + // nested compute region. + case OpenACCGangKind::Num: + // This isn't allowed if there is a 'num_gangs' on the kernel construct, + // and makes loop-with-gang-clause ill-formed inside of this 'loop', but + // nothing can be enforced here. + return ExprResult{E}; + case OpenACCGangKind::Static: + return CheckGangStaticExpr(*this, E); + } + } break; + case OpenACCDirectiveKind::Serial: { + switch (GK) { + // 'dim' and 'num' don't really make sense on serial, and GCC rejects them + // too, so we disallow them too. + case OpenACCGangKind::Dim: + case OpenACCGangKind::Num: + Diag(E->getBeginLoc(), diag::err_acc_gang_arg_invalid) + << GK << /*Kernels=*/3; + return ExprError(); + case OpenACCGangKind::Static: + return CheckGangStaticExpr(*this, E); + } + } + default: + llvm_unreachable("Non compute construct in active compute construct?"); + } + + llvm_unreachable("Compute construct directive not handled?"); +} + ExprResult SemaOpenACC::CheckTileSizeExpr(Expr *SizeExpr) { if (!SizeExpr) return ExprError(); @@ -2031,7 +2259,7 @@ StmtResult SemaOpenACC::ActOnEndStmtDirective(OpenACCDirectiveKind K, // If we are in the scope of a compute construct, add this to the list of // loop constructs that need assigning to the next closing compute // construct. - if (InsideComputeConstruct) + if (isInComputeConstruct()) ParentlessLoopConstructs.push_back(LoopConstruct); return LoopConstruct; diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h index 5753c9eccf6c92..cde40773336866 100644 --- a/clang/lib/Sema/TreeTransform.h +++ b/clang/lib/Sema/TreeTransform.h @@ -11940,6 +11940,29 @@ void OpenACCClauseTransform<Derived>::VisitTileClause( ParsedClause.getLParenLoc(), ParsedClause.getIntExprs(), ParsedClause.getEndLoc()); } +template <typename Derived> +void OpenACCClauseTransform<Derived>::VisitGangClause( + const OpenACCGangClause &C) { + llvm::SmallVector<OpenACCGangKind> TransformedGangKinds; + llvm::SmallVector<Expr *> TransformedIntExprs; + + for (unsigned I = 0; I < C.getNumExprs(); ++I) { + ExprResult ER = Self.TransformExpr(const_cast<Expr *>(C.getExpr(I).second)); + if (!ER.isUsable()) + continue; + + ER = Self.getSema().OpenACC().CheckGangExpr(C.getExpr(I).first, ER.get()); + if (!ER.isUsable()) + continue; + TransformedGangKinds.push_back(C.getExpr(I).first); + TransformedIntExprs.push_back(ER.get()); + } + + NewClause = OpenACCGangClause::Create( + Self.getSema().getASTContext(), ParsedClause.getBeginLoc(), + ParsedClause.getLParenLoc(), TransformedGangKinds, TransformedIntExprs, + ParsedClause.getEndLoc()); +} } // namespace template <typename Derived> OpenACCClause *TreeTransform<Derived>::TransformOpenACCClause( diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp index e638129897692f..0339419da43cab 100644 --- a/clang/lib/Serialization/ASTReader.cpp +++ b/clang/lib/Serialization/ASTReader.cpp @@ -12326,6 +12326,18 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() { return OpenACCTileClause::Create(getContext(), BeginLoc, LParenLoc, SizeExprs, EndLoc); } + case OpenACCClauseKind::Gang: { + SourceLocation LParenLoc = readSourceLocation(); + unsigned NumExprs = readInt(); + llvm::SmallVector<OpenACCGangKind> GangKinds; + llvm::SmallVector<Expr *> Exprs; + for (unsigned I = 0; I < NumExprs; ++I) { + GangKinds.push_back(readEnum<OpenACCGangKind>()); + Exprs.push_back(readSubExpr()); + } + return OpenACCGangClause::Create(getContext(), BeginLoc, LParenLoc, + GangKinds, Exprs, EndLoc); + } case OpenACCClauseKind::Finalize: case OpenACCClauseKind::IfPresent: @@ -12342,7 +12354,6 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() { case OpenACCClauseKind::Bind: case OpenACCClauseKind::DeviceNum: case OpenACCClauseKind::DefaultAsync: - case OpenACCClauseKind::Gang: 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 4976327fc654ee..583d9a4bccb800 100644 --- a/clang/lib/Serialization/ASTWriter.cpp +++ b/clang/lib/Serialization/ASTWriter.cpp @@ -8182,6 +8182,16 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) { AddStmt(E); return; } + case OpenACCClauseKind::Gang: { + const auto *GC = cast<OpenACCGangClause>(C); + writeSourceLocation(GC->getLParenLoc()); + writeUInt32(GC->getNumExprs()); + for (unsigned I = 0; I < GC->getNumExprs(); ++I) { + writeEnum(GC->getExpr(I).first); + AddStmt(const_cast<Expr *>(GC->getExpr(I).second)); + } + return; + } case OpenACCClauseKind::Finalize: case OpenACCClauseKind::IfPresent: @@ -8198,7 +8208,6 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) { case OpenACCClauseKind::Bind: case OpenACCClauseKind::DeviceNum: case OpenACCClauseKind::DefaultAsync: - case OpenACCClauseKind::Gang: case OpenACCClauseKind::Invalid: llvm_unreachable("Clause serialization not yet implemented"); } diff --git a/clang/test/AST/ast-print-openacc-loop-construct.cpp b/clang/test/AST/ast-print-openacc-loop-construct.cpp index aee4591cab428f..baa4b173f88edc 100644 --- a/clang/test/AST/ast-print-openacc-loop-construct.cpp +++ b/clang/test/AST/ast-print-openacc-loop-construct.cpp @@ -95,4 +95,86 @@ void foo() { for(;;) for(;;) for(;;); + +// CHECK: #pragma acc loop gang(dim: 2) +// CHECK-NEXT: for (;;) +// CHECK-NEXT: ; +#pragma acc loop gang(dim:2) + for(;;); + +// CHECK: #pragma acc loop gang(static: i) +// CHECK-NEXT: for (;;) +// CHECK-NEXT: ; +#pragma acc loop gang(static:i) + for(;;); + +// CHECK: #pragma acc loop gang(static: i) gang(dim: 2) +// CHECK-NEXT: for (;;) +// CHECK-NEXT: ; +#pragma acc loop gang(static:i) gang(dim:2) + for(;;); + +// CHECK: #pragma acc parallel +// CHECK-NEXT: #pragma acc loop gang(dim: 2) +// CHECK-NEXT: for (;;) +// CHECK-NEXT: ; +#pragma acc parallel +#pragma acc loop gang(dim:2) + for(;;); + +// CHECK: #pragma acc parallel +// CHECK-NEXT: #pragma acc loop gang(static: i) +// CHECK-NEXT: for (;;) +// CHECK-NEXT: ; +#pragma acc parallel +#pragma acc loop gang(static:i) + for(;;); + +// CHECK: #pragma acc parallel +// CHECK-NEXT: #pragma acc loop gang(static: i) gang(dim: 2) +// CHECK-NEXT: for (;;) +// CHECK-NEXT: ; +#pragma acc parallel +#pragma acc loop gang(static:i) gang(dim:2) + for(;;); + +// CHECK: #pragma acc kernels +// CHECK-NEXT: #pragma acc loop gang(num: i) gang(static: i) +// CHECK-NEXT: for (;;) +// CHECK-NEXT: ; +#pragma acc kernels +#pragma acc loop gang(i) gang(static:i) + for(;;); + +// CHECK: #pragma acc kernels +// CHECK-NEXT: #pragma acc loop gang(num: i) gang(static: i) +// CHECK-NEXT: for (;;) +// CHECK-NEXT: ; +#pragma acc kernels +#pragma acc loop gang(num:i) gang(static:i) + for(;;); + +// CHECK: #pragma acc serial +// CHECK-NEXT: #pragma acc loop gang(static: i) +// CHECK-NEXT: for (;;) +// CHECK-NEXT: ; +#pragma acc serial +#pragma acc loop gang(static:i) + for(;;); + +// CHECK: #pragma acc serial +// CHECK-NEXT: #pragma acc loop gang(static: *) +// CHECK-NEXT: for (;;) +// CHECK-NEXT: ; +#pragma acc serial +#pragma acc loop gang(static:*) + for(;;); + +// CHECK: #pragma acc serial +// CHECK-NEXT: #pragma acc loop +// CHECK-NEXT: for (;;) +// CHECK-NEXT: ; +#pragma acc serial +#pragma acc loop gang + for(;;); } diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c index 6c382379a8a7ea..899fbd78b87298 100644 --- a/clang/test/ParserOpenACC/parse-clauses.c +++ b/clang/test/ParserOpenACC/parse-clauses.c @@ -1202,7 +1202,6 @@ void Tile() { } void Gang() { - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}} #pragma acc loop gang for(;;){} // expected-error@+3{{expected expression}} @@ -1210,68 +1209,58 @@ void Gang() { // expected-note@+1{{to match this '('}} #pragma acc loop gang( for(;;){} - // expected-error@+2{{expected expression}} - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}} + // expected-error@+1{{expected expression}} #pragma acc loop gang() for(;;){} - // expected-error@+2{{expected expression}} - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}} + // expected-error@+1{{expected expression}} #pragma acc loop gang(5, *) for(;;){} - // expected-error@+2{{expected expression}} - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}} + // expected-error@+1{{expected expression}} #pragma acc loop gang(*) for(;;){} - // expected-error@+2{{expected expression}} - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}} + // expected-error@+1{{expected expression}} #pragma acc loop gang(5, num:*) for(;;){} - // expected-error@+2{{expected expression}} - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}} + // expected-error@+1{{expected expression}} #pragma acc loop gang(num:5, *) for(;;){} - // expected-error@+2{{expected expression}} - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}} + // expected-error@+1{{expected expression}} #pragma acc loop gang(num:5, num:*) for(;;){} - // expected-error@+2{{expected expression}} - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}} + // expected-error@+1{{expected expression}} #pragma acc loop gang(num:*) for(;;){} - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}} -#pragma acc loop gang(dim:5) +#pragma acc loop gang(dim:2) for(;;){} - // expected-error@+2{{expected expression}} - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}} + // expected-error@+1{{expected expression}} #pragma acc loop gang(dim:5, dim:*) for(;;){} - // expected-error@+2{{expected expression}} - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}} + // expected-error@+1{{expected expression}} #pragma acc loop gang(dim:*) for(;;){} - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}} #pragma acc loop gang(static:*) for(;;){} - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}} + // expected-error@+2{{OpenACC 'gang' clause may have at most one 'static' argument}} + // expected-note@+1{{previous expression is here}} #pragma acc loop gang(static:*, static:5) for(;;){} - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}} +#pragma acc kernels #pragma acc loop gang(static:*, 5) for(;;){} - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}} +#pragma acc kernels #pragma acc loop gang(static:45, 5) for(;;){} @@ -1330,11 +1319,16 @@ void Gang() { #pragma acc loop gang(dim:45 for(;;){} - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}} -#pragma acc loop gang(static:*, dim:returns_int(), 5) +#pragma acc kernels +#pragma acc loop gang(static:*, 5) + for(;;){} + + // expected-error@+1{{argument to 'gang' clause dimension must be a constant expression}} +#pragma acc loop gang(static:*, dim:returns_int()) for(;;){} - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}} + // expected-error@+2 2{{'num' argument on 'gang' clause is not permitted on an orphaned 'loop' construct}} + // expected-error@+1{{argument to 'gang' clause dimension must be a constant expression}} #pragma acc loop gang(num: 32, static:*, dim:returns_int(), 5) for(;;){} diff --git a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c index d08497a7782edb..89000517c43fb5 100644 --- a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c +++ b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c @@ -214,8 +214,7 @@ void uses() { // expected-error@+1{{OpenACC 'tile' clause is not valid on 'kernels' directive}} #pragma acc kernels device_type(*) tile(Var, 1) while(1); - // expected-error@+2{{OpenACC clause 'gang' may not follow a 'dtype' clause in a compute construct}} - // expected-note@+1{{previous clause is here}} + // expected-error@+1{{OpenACC 'gang' clause is not valid on 'kernels' directive}} #pragma acc kernels dtype(*) gang while(1); #pragma acc kernels device_type(*) wait 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 3da7f0e9836be8..6c2c79b02a4131 100644 --- a/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c +++ b/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c @@ -168,7 +168,6 @@ void uses() { #pragma acc loop auto tile(1+2, 1) for(;;) for(;;); - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented}} #pragma acc loop auto gang for(;;); // expected-error@+1{{OpenACC 'wait' clause is not valid on 'loop' directive}} @@ -306,7 +305,6 @@ void uses() { #pragma acc loop tile(1+2, 1) auto for(;;) for(;;); - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented}} #pragma acc loop gang auto for(;;); // expected-error@+1{{OpenACC 'wait' clause is not valid on 'loop' directive}} @@ -445,7 +443,6 @@ void uses() { #pragma acc loop independent tile(1+2, 1) for(;;) for(;;); - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented}} #pragma acc loop independent gang for(;;); // expected-error@+1{{OpenACC 'wait' clause is not valid on 'loop' directive}} @@ -583,7 +580,6 @@ void uses() { #pragma acc loop tile(1+2, 1) independent for(;;) for(;;); - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented}} #pragma acc loop gang independent for(;;); // expected-error@+1{{OpenACC 'wait' clause is not valid on 'loop' directive}} @@ -591,9 +587,8 @@ void uses() { for(;;); // 'seq' cannot be combined with 'gang', 'worker' or 'vector' - // expected-error@+3{{OpenACC clause 'gang' may not appear on the same construct as a 'seq' clause on a 'loop' construct}} - // expected-note@+2{{previous clause is here}} - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented}} + // 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(;;); // expected-error@+3{{OpenACC clause 'worker' may not appear on the same construct as a 'seq' clause on a 'loop' construct}} @@ -735,10 +730,8 @@ void uses() { #pragma acc loop seq wait for(;;); - // TODO OpenACC: when 'gang' is implemented and makes it to the AST, this should diagnose because of a conflict with 'seq'. - // TODOexpected-error@+3{{OpenACC clause 'gang' may not appear on the same construct as a 'seq' clause on a 'loop' construct}} - // TODOexpected-note@+2{{previous clause is here}} - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented}} + // 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(;;); // TODO OpenACC: when 'worker' is implemented and makes it to the AST, this should diagnose because of a conflict with 'seq'. diff --git a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c index 3d77c031f42630..cedef3ca858f5e 100644 --- a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c +++ b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c @@ -193,7 +193,6 @@ void uses() { for(;;) for(;;); - // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}} #pragma acc loop dtype(*) gang for(;;); // expected-error@+1{{OpenACC 'wait' clause is not valid on 'loop' directive}} diff --git a/clang/test/SemaOpenACC/loop-construct-gang-ast.cpp b/clang/test/SemaOpenACC/loop-construct-gang-ast.cpp new file mode 100644 index 00000000000000..e797d842e240dc --- /dev/null +++ b/clang/test/SemaOpenACC/loop-construct-gang-ast.cpp @@ -0,0 +1,330 @@ +// RUN: %clang_cc1 %s -fopenacc -ast-dump | FileCheck %s + +// Test this with PCH. +// RUN: %clang_cc1 %s -fopenacc -emit-pch -o %t %s +// RUN: %clang_cc1 %s -fopenacc -include-pch %t -ast-dump-all | FileCheck %s +#ifndef PCH_HELPER +#define PCH_HELPER + +void NormalUses() { + // CHECK: FunctionDecl{{.*}}NormalUses + // CHECK-NEXT: CompoundStmt + + // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan> + // CHECK-NEXT: gang clause dim + // CHECK-NEXT: ConstantExpr{{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}}'int' 1 + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: NullStmt +#pragma acc loop gang(dim:1) + for(;;); + + int Val; + // CHECK-NEXT: DeclStmt + // CHECK-NEXT: VarDecl{{.*}} used Val 'int' + + // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan> + // CHECK-NEXT: gang clause static + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue> + // CHECK-NEXT: DeclRefExpr{{.*}} 'Val' 'int' + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: NullStmt +#pragma acc loop gang(static:Val) + for(;;); + + // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} kernels + // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]] + // CHECK-NEXT: gang clause num + // CHECK-NEXT: IntegerLiteral{{.*}}'int' 1 + // CHECK-NEXT: gang clause static + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue> + // CHECK-NEXT: DeclRefExpr{{.*}} 'Val' 'int' + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: NullStmt +#pragma acc kernels +#pragma acc loop gang(num:1) gang(static:Val) + for(;;); + + // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} parallel + // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]] + // CHECK-NEXT: gang clause dim static + // CHECK-NEXT: ConstantExpr{{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}}'int' 1 + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue> + // CHECK-NEXT: DeclRefExpr{{.*}} 'Val' 'int' + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: NullStmt +#pragma acc parallel +#pragma acc loop gang(dim:1, static:Val) + for(;;); + + // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} serial + // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]] + // CHECK-NEXT: gang clause static + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue> + // CHECK-NEXT: DeclRefExpr{{.*}} 'Val' 'int' + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: NullStmt +#pragma acc serial +#pragma acc loop gang(static:Val) + for(;;); + + // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} serial + // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]] + // CHECK-NEXT: gang clause static + // CHECK-NEXT: OpenACCAsteriskSizeExpr + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: NullStmt +#pragma acc serial +#pragma acc loop gang(static:*) + for(;;); + + // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} serial + // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]] + // CHECK-NEXT: gang clause + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: NullStmt +#pragma acc serial +#pragma acc loop gang + for(;;); +} + +template<typename T, unsigned One> +void TemplateUses(T Val) { + // CHECK: FunctionTemplateDecl{{.*}}TemplateUses + // CHECK-NEXT: TemplateTypeParmDecl {{.*}} referenced typename depth 0 index 0 T + // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} referenced 'unsigned int' depth 0 index 1 One + // CHECK-NEXT: FunctionDecl{{.*}} TemplateUses 'void (T)' + // CHECK-NEXT: ParmVarDecl{{.*}} referenced Val 'T' + // CHECK-NEXT: CompoundStmt + + // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan> + // CHECK-NEXT: gang clause dim + // CHECK-NEXT: DeclRefExpr{{.*}}'unsigned int' NonTypeTemplateParm{{.*}} 'One' 'unsigned int' + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: NullStmt +#pragma acc loop gang(dim:One) + for(;;); + + // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan> + // CHECK-NEXT: gang clause static + // CHECK-NEXT: DeclRefExpr{{.*}}'T' lvalue ParmVar{{.*}} 'Val' 'T' + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: NullStmt +#pragma acc loop gang(static:Val) + for(;;); + + // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan> + // CHECK-NEXT: gang clause static + // CHECK-NEXT: OpenACCAsteriskSizeExpr + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: NullStmt +#pragma acc loop gang(static:*) + for(;;); + + // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} parallel + // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]] + // CHECK-NEXT: gang clause dim + // CHECK-NEXT: DeclRefExpr{{.*}}'unsigned int' NonTypeTemplateParm{{.*}} 'One' 'unsigned int' + // CHECK-NEXT: gang clause static + // CHECK-NEXT: DeclRefExpr{{.*}}'T' lvalue ParmVar{{.*}} 'Val' 'T' + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: NullStmt +#pragma acc parallel +#pragma acc loop gang(dim:One) gang(static:Val) + for(;;); + + // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} parallel + // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]] + // CHECK-NEXT: gang clause dim static + // CHECK-NEXT: DeclRefExpr{{.*}}'unsigned int' NonTypeTemplateParm{{.*}} 'One' 'unsigned int' + // CHECK-NEXT: DeclRefExpr{{.*}}'T' lvalue ParmVar{{.*}} 'Val' 'T' + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: NullStmt +#pragma acc parallel +#pragma acc loop gang(dim:One, static:Val) + for(;;); + + // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} serial + // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]] + // CHECK-NEXT: gang clause static + // CHECK-NEXT: DeclRefExpr{{.*}}'T' lvalue ParmVar{{.*}} 'Val' 'T' + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: NullStmt +#pragma acc serial +#pragma acc loop gang(static:Val) + for(;;); + + // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} serial + // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]] + // CHECK-NEXT: gang clause + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: NullStmt +#pragma acc serial +#pragma acc loop gang + for(;;); + + // Instantiation: + // CHECK-NEXT: FunctionDecl{{.*}} used TemplateUses 'void (int)' implicit_instantiation + // CHECK-NEXT: TemplateArgument type 'int' + // CHECK-NEXT: BuiltinType{{.*}} 'int' + // CHECK-NEXT: TemplateArgument integral '1U' + // CHECK-NEXT: ParmVarDecl{{.*}} used Val 'int' + // CHECK-NEXT: CompoundStmt + // + // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan> + // CHECK-NEXT: gang clause dim + // CHECK-NEXT: ConstantExpr{{.*}} 'unsigned int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}}'unsigned int' + // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'unsigned int' depth 0 index 1 One + // CHECK-NEXT: IntegerLiteral{{.*}}'unsigned int' 1 + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: NullStmt + // + // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan> + // CHECK-NEXT: gang clause static + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue> + // CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue ParmVar{{.*}} 'Val' 'int' + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: NullStmt + // + // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan> + // CHECK-NEXT: gang clause static + // CHECK-NEXT: OpenACCAsteriskSizeExpr + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: NullStmt + // + // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} parallel + // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]] + // CHECK-NEXT: gang clause dim + // CHECK-NEXT: ConstantExpr{{.*}} 'unsigned int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}}'unsigned int' + // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'unsigned int' depth 0 index 1 One + // CHECK-NEXT: IntegerLiteral{{.*}}'unsigned int' 1 + // CHECK-NEXT: gang clause static + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue> + // CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue ParmVar{{.*}} 'Val' 'int' + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: NullStmt + // + // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} parallel + // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]] + // CHECK-NEXT: gang clause dim static + // CHECK-NEXT: ConstantExpr{{.*}} 'unsigned int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}}'unsigned int' + // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'unsigned int' depth 0 index 1 One + // CHECK-NEXT: IntegerLiteral{{.*}}'unsigned int' 1 + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue> + // CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue ParmVar{{.*}} 'Val' 'int' + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: NullStmt + // + // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} serial + // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]] + // CHECK-NEXT: gang clause static + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue> + // CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue ParmVar{{.*}} 'Val' 'int' + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: NullStmt + // + // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} serial + // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]] + // CHECK-NEXT: gang clause + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: <<<NULL>> + // CHECK-NEXT: NullStmt +} + +void inst() { + TemplateUses<int, 1>(5); +} + +#endif // PCH_HELPER diff --git a/clang/test/SemaOpenACC/loop-construct-gang-clause.cpp b/clang/test/SemaOpenACC/loop-construct-gang-clause.cpp new file mode 100644 index 00000000000000..ab6439ae576193 --- /dev/null +++ b/clang/test/SemaOpenACC/loop-construct-gang-clause.cpp @@ -0,0 +1,335 @@ +// RUN: %clang_cc1 %s -fopenacc -verify + +struct S{}; +struct Converts{ + operator int(); +}; + +template<typename T, unsigned Zero, unsigned Two, unsigned Four> +void ParallelOrOrphanTempl() { + T i; + // expected-error@+1{{'num' argument on 'gang' clause is not permitted on an orphaned 'loop' construct}} +#pragma acc loop gang(i) + for(;;); + // expected-error@+1{{'num' argument on 'gang' clause is not permitted on an orphaned 'loop' construct}} +#pragma acc loop gang(num:i) + for(;;); + + // expected-error@+2{{'num' argument on 'gang' clause is not permitted on a 'loop' construct associated with a 'parallel' compute construct}} +#pragma acc parallel +#pragma acc loop gang(i) + for(;;); + + // expected-error@+2{{'num' argument on 'gang' clause is not permitted on a 'loop' construct associated with a 'parallel' compute construct}} +#pragma acc parallel +#pragma acc loop gang(num:i) + for(;;); + + // expected-error@+1{{argument to 'gang' clause dimension must be a constant expression}} +#pragma acc loop gang(dim:i) + for(;;); + + // expected-error@+2{{argument to 'gang' clause dimension must be a constant expression}} +#pragma acc parallel +#pragma acc loop gang(dim:i) + for(;;); + + // expected-error@+1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 0}} +#pragma acc loop gang(dim:Zero) + for(;;); + + // expected-error@+2{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 0}} +#pragma acc parallel +#pragma acc loop gang(dim:Zero) + for(;;); + + // expected-error@+1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 4}} +#pragma acc loop gang(dim:Four) + for(;;); + + // expected-error@+2{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 4}} +#pragma acc parallel +#pragma acc loop gang(dim:Four) + for(;;); + +#pragma acc loop gang(static:i) gang(dim:Two) + for(;;); + +#pragma acc parallel +#pragma acc loop gang(dim:Two) gang(static:*) + for(;;); + +#pragma acc parallel +#pragma acc loop gang(dim:Two, static:i) + for(;;); + + // expected-error@+4{{OpenACC 'gang' clause may have at most one 'static' argument}} + // expected-note@+3{{previous expression is here}} + // expected-error@+2{{OpenACC 'gang' clause may have at most one 'dim' argument}} + // expected-note@+1{{previous expression is here}} +#pragma acc loop gang(static:i, static:i, dim:Two, dim:1) + for(;;); +} + +void ParallelOrOrphan() { + ParallelOrOrphanTempl<int, 0, 2, 4>(); // expected-note{{in instantiation of function template}} + + int i; + // expected-error@+1{{'num' argument on 'gang' clause is not permitted on an orphaned 'loop' construct}} +#pragma acc loop gang(i) + for(;;); + // expected-error@+1{{'num' argument on 'gang' clause is not permitted on an orphaned 'loop' construct}} +#pragma acc loop gang(num:i) + for(;;); + + // expected-error@+2{{'num' argument on 'gang' clause is not permitted on a 'loop' construct associated with a 'parallel' compute construct}} +#pragma acc parallel +#pragma acc loop gang(i) + for(;;); + + // expected-error@+2{{'num' argument on 'gang' clause is not permitted on a 'loop' construct associated with a 'parallel' compute construct}} +#pragma acc parallel +#pragma acc loop gang(num:i) + for(;;); + + // expected-error@+1{{argument to 'gang' clause dimension must be a constant expression}} +#pragma acc loop gang(dim:i) + for(;;); + + // expected-error@+2{{argument to 'gang' clause dimension must be a constant expression}} +#pragma acc parallel +#pragma acc loop gang(dim:i) + for(;;); + + // expected-error@+1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 0}} +#pragma acc loop gang(dim:0) + for(;;); + + // expected-error@+2{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 0}} +#pragma acc parallel +#pragma acc loop gang(dim:0) + for(;;); + + // expected-error@+1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 4}} +#pragma acc loop gang(dim:4) + for(;;); + + // expected-error@+2{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 4}} +#pragma acc parallel +#pragma acc loop gang(dim:4) + for(;;); + +#pragma acc loop gang(static:i) gang(dim:2) + for(;;); + +#pragma acc parallel +#pragma acc loop gang(dim:2) gang(static:i) + for(;;); + + S s; + // expected-error@+2{{OpenACC clause 'gang' requires expression of integer type ('S' invalid)}} +#pragma acc parallel +#pragma acc loop gang(dim:2) gang(static:s) + for(;;); + + Converts C; +#pragma acc parallel +#pragma acc loop gang(dim:2) gang(static:C) + for(;;); +} + +template<typename SomeS, typename SomeC, typename Int> +void StaticIsIntegralTempl() { + SomeS s; + // expected-error@+2{{OpenACC clause 'gang' requires expression of integer type ('S' invalid)}} +#pragma acc parallel +#pragma acc loop gang(dim:2) gang(static:s) + for(;;); + + SomeC C; +#pragma acc parallel +#pragma acc loop gang(dim:2) gang(static:C) + for(;;); + Int i; +#pragma acc parallel +#pragma acc loop gang(dim:2) gang(static:i) + for(;;); + +#pragma acc parallel +#pragma acc loop gang(dim:2) gang(static:*) + for(;;); +} + +void StaticIsIntegral() { + StaticIsIntegralTempl<S, Converts, int>();// expected-note{{in instantiation of function template}} + + S s; + // expected-error@+2{{OpenACC clause 'gang' requires expression of integer type ('S' invalid)}} +#pragma acc parallel +#pragma acc loop gang(dim:2) gang(static:s) + for(;;); + + Converts C; +#pragma acc parallel +#pragma acc loop gang(dim:2) gang(static:C) + for(;;); +} + +template<unsigned I> +void SerialTempl() { + // expected-error@+2{{'num' argument on 'gang' clause is not permitted on a 'loop' construct associated with a 'serial' compute construct}} +#pragma acc serial +#pragma acc loop gang(I) + for(;;); + + // expected-error@+2{{'num' argument on 'gang' clause is not permitted on a 'loop' construct associated with a 'serial' compute construct}} +#pragma acc serial +#pragma acc loop gang(num:I) + for(;;); + + // expected-error@+2{{'dim' argument on 'gang' clause is not permitted on a 'loop' construct associated with a 'serial' compute construct}} +#pragma acc serial +#pragma acc loop gang(dim:I) + for(;;); + +#pragma acc serial +#pragma acc loop gang(static:I) + for(;;); +} + +void Serial() { + SerialTempl<2>(); + + // expected-error@+2{{'num' argument on 'gang' clause is not permitted on a 'loop' construct associated with a 'serial' compute construct}} +#pragma acc serial +#pragma acc loop gang(1) + for(;;); + + // expected-error@+2{{'num' argument on 'gang' clause is not permitted on a 'loop' construct associated with a 'serial' compute construct}} +#pragma acc serial +#pragma acc loop gang(num:1) + for(;;); + + // expected-error@+2{{'dim' argument on 'gang' clause is not permitted on a 'loop' construct associated with a 'serial' compute construct}} +#pragma acc serial +#pragma acc loop gang(dim:1) + for(;;); + +#pragma acc serial +#pragma acc loop gang(static:1) + for(;;); + + int i; + +#pragma acc serial +#pragma acc loop gang(static:i) + for(;;); +} + +template<typename T> +void KernelsTempl() { + T t; + // expected-error@+2{{'dim' argument on 'gang' clause is not permitted on a 'loop' construct associated with a 'kernels' compute construct}} +#pragma acc kernels +#pragma acc loop gang(dim:t) + for(;;); + +#pragma acc kernels +#pragma acc loop gang(static:t) + for(;;); + + // expected-error@+3{{'num' argument to 'gang' clause not allowed on a 'loop' construct associated with a 'kernels' construct that has a 'num_gangs' clause}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels num_gangs(t) +#pragma acc loop gang(t) + for(;;); + + // expected-error@+3{{'num' argument to 'gang' clause not allowed on a 'loop' construct associated with a 'kernels' construct that has a 'num_gangs' clause}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels num_gangs(t) +#pragma acc loop gang(num:t) + for(;;); +} + +void Kernels() { + KernelsTempl<unsigned>(); + + // expected-error@+2{{'dim' argument on 'gang' clause is not permitted on a 'loop' construct associated with a 'kernels' compute construct}} +#pragma acc kernels +#pragma acc loop gang(dim:1) + for(;;); + unsigned t; +#pragma acc kernels +#pragma acc loop gang(static:t) + for(;;); + + // expected-error@+3{{'num' argument to 'gang' clause not allowed on a 'loop' construct associated with a 'kernels' construct that has a 'num_gangs' clause}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels num_gangs(1) +#pragma acc loop gang(1) + for(;;); + + // expected-error@+3{{'num' argument to 'gang' clause not allowed on a 'loop' construct associated with a 'kernels' construct that has a 'num_gangs' clause}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels num_gangs(1) +#pragma acc loop gang(num:1) + for(;;); + +#pragma acc kernels +#pragma acc loop gang(num:1) + for(;;) { + // expected-error@+2{{loop with a 'gang' clause may not exist in the region of a 'gang' clause on a 'kernels' compute construct}} + // expected-note@-3{{previous clause is here}} +#pragma acc loop gang(static:1) + for(;;); + } + +#pragma acc kernels +#pragma acc loop gang(num:1) + for(;;) { + // allowed, intervening compute construct +#pragma acc serial +#pragma acc loop gang(static:1) + for(;;); + } + +#pragma acc kernels +#pragma acc loop gang(num:1) + for(;;); + + // OK, on a diff erent 'loop', not in the assoc statement. +#pragma acc loop gang(static:1) + for(;;); + + // expected-error@+3{{OpenACC 'gang' clause may have at most one unnamed or 'num' argument}} + // expected-note@+2{{previous expression is here}} +#pragma acc kernels +#pragma acc loop gang(5, num:1) + for(;;); + + // expected-error@+3{{OpenACC 'gang' clause may have at most one unnamed or 'num' argument}} + // expected-note@+2{{previous expression is here}} +#pragma acc kernels +#pragma acc loop gang(num:5, 1) + for(;;); + + // expected-error@+3{{OpenACC 'gang' clause may have at most one unnamed or 'num' argument}} + // expected-note@+2{{previous expression is here}} +#pragma acc kernels +#pragma acc loop gang(num:5, num:1) + for(;;); +} + +void MaxOneEntry() { + // expected-error@+3{{OpenACC 'gang' clause may have at most one 'static' argument}} + // expected-note@+2{{previous expression is here}} +#pragma acc kernels +#pragma acc loop gang(static: 1, static:1) + for(;;); + +#pragma acc kernels +#pragma acc loop gang gang(static:1) + for(;;); +} + + diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp index c282a9071391e7..2ffe47fbd74476 100644 --- a/clang/tools/libclang/CIndex.cpp +++ b/clang/tools/libclang/CIndex.cpp @@ -2907,6 +2907,11 @@ void OpenACCClauseEnqueue::VisitSeqClause(const OpenACCSeqClause &C) {} void OpenACCClauseEnqueue::VisitCollapseClause(const OpenACCCollapseClause &C) { Visitor.AddStmt(C.getLoopCount()); } +void OpenACCClauseEnqueue::VisitGangClause(const OpenACCGangClause &C) { + for (unsigned I = 0; I < C.getNumExprs(); ++I) { + Visitor.AddStmt(C.getExpr(I).second); + } +} } // namespace void EnqueueVisitor::EnqueueChildren(const OpenACCClause *C) { _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
