Author: Erich Keane Date: 2024-04-22T08:57:25-07:00 New Revision: dc20a0ea1fd04a2ef95eb5c73e9f88357fc5f694
URL: https://github.com/llvm/llvm-project/commit/dc20a0ea1fd04a2ef95eb5c73e9f88357fc5f694 DIFF: https://github.com/llvm/llvm-project/commit/dc20a0ea1fd04a2ef95eb5c73e9f88357fc5f694.diff LOG: [OpenACC] Implement 'num_gangs' sema for compute constructs (#89460) num_gangs takes an 'int-expr-list', for 'parallel', and an 'int-expr' for 'kernels'. This patch changes the parsing to always parse it as an 'int-expr-list', then correct the expression count during Sema. It also implements the rest of the semantic analysis changes for this clause. Added: clang/test/SemaOpenACC/compute-construct-num_gangs-clause.c clang/test/SemaOpenACC/compute-construct-num_gangs-clause.cpp Modified: clang/include/clang/AST/OpenACCClause.h clang/include/clang/Basic/DiagnosticSemaKinds.td clang/include/clang/Basic/OpenACCClauses.def 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/ParserOpenACC/parse-clauses.c clang/test/SemaOpenACC/compute-construct-intexpr-clause-ast.cpp clang/tools/libclang/CIndex.cpp Removed: ################################################################################ diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h index 8b6d3221aa066b..277a351c49fcb8 100644 --- a/clang/include/clang/AST/OpenACCClause.h +++ b/clang/include/clang/AST/OpenACCClause.h @@ -156,33 +156,88 @@ class OpenACCSelfClause : public OpenACCClauseWithCondition { Expr *ConditionExpr, SourceLocation EndLoc); }; -/// Represents oen of a handful of classes that have a single integer +/// Represents a clause that has one or more IntExprs. It does not own the +/// IntExprs, but provides 'children' and other accessors. +class OpenACCClauseWithIntExprs : public OpenACCClauseWithParams { + MutableArrayRef<Expr *> IntExprs; + +protected: + OpenACCClauseWithIntExprs(OpenACCClauseKind K, SourceLocation BeginLoc, + SourceLocation LParenLoc, SourceLocation EndLoc) + : OpenACCClauseWithParams(K, BeginLoc, LParenLoc, EndLoc) {} + + /// Used only for initialization, the leaf class can initialize this to + /// trailing storage. + void setIntExprs(MutableArrayRef<Expr *> NewIntExprs) { + assert(IntExprs.empty() && "Cannot change IntExprs list"); + IntExprs = NewIntExprs; + } + + /// Gets the entire list of integer expressions, but leave it to the + /// individual clauses to expose this how they'd like. + llvm::ArrayRef<Expr *> getIntExprs() const { return IntExprs; } + +public: + child_range children() { + return child_range(reinterpret_cast<Stmt **>(IntExprs.begin()), + reinterpret_cast<Stmt **>(IntExprs.end())); + } + + const_child_range children() const { + child_range Children = + const_cast<OpenACCClauseWithIntExprs *>(this)->children(); + return const_child_range(Children.begin(), Children.end()); + } +}; + +class OpenACCNumGangsClause final + : public OpenACCClauseWithIntExprs, + public llvm::TrailingObjects<OpenACCNumGangsClause, Expr *> { + + OpenACCNumGangsClause(SourceLocation BeginLoc, SourceLocation LParenLoc, + ArrayRef<Expr *> IntExprs, SourceLocation EndLoc) + : OpenACCClauseWithIntExprs(OpenACCClauseKind::NumGangs, BeginLoc, + LParenLoc, EndLoc) { + std::uninitialized_copy(IntExprs.begin(), IntExprs.end(), + getTrailingObjects<Expr *>()); + setIntExprs(MutableArrayRef(getTrailingObjects<Expr *>(), IntExprs.size())); + } + +public: + static OpenACCNumGangsClause * + Create(const ASTContext &C, SourceLocation BeginLoc, SourceLocation LParenLoc, + ArrayRef<Expr *> IntExprs, SourceLocation EndLoc); + + llvm::ArrayRef<Expr *> getIntExprs() { + return OpenACCClauseWithIntExprs::getIntExprs(); + } + + llvm::ArrayRef<Expr *> getIntExprs() const { + return OpenACCClauseWithIntExprs::getIntExprs(); + } +}; + +/// Represents one of a handful of clauses that have a single integer /// expression. -class OpenACCClauseWithSingleIntExpr : public OpenACCClauseWithParams { +class OpenACCClauseWithSingleIntExpr : public OpenACCClauseWithIntExprs { Expr *IntExpr; protected: OpenACCClauseWithSingleIntExpr(OpenACCClauseKind K, SourceLocation BeginLoc, SourceLocation LParenLoc, Expr *IntExpr, SourceLocation EndLoc) - : OpenACCClauseWithParams(K, BeginLoc, LParenLoc, EndLoc), - IntExpr(IntExpr) {} + : OpenACCClauseWithIntExprs(K, BeginLoc, LParenLoc, EndLoc), + IntExpr(IntExpr) { + setIntExprs(MutableArrayRef<Expr *>{&this->IntExpr, 1}); + } public: - bool hasIntExpr() const { return IntExpr; } - const Expr *getIntExpr() const { return IntExpr; } - - Expr *getIntExpr() { return IntExpr; }; - - child_range children() { - return child_range(reinterpret_cast<Stmt **>(&IntExpr), - reinterpret_cast<Stmt **>(&IntExpr + 1)); + bool hasIntExpr() const { return !getIntExprs().empty(); } + const Expr *getIntExpr() const { + return hasIntExpr() ? getIntExprs()[0] : nullptr; } - const_child_range children() const { - return const_child_range(reinterpret_cast<Stmt *const *>(&IntExpr), - reinterpret_cast<Stmt *const *>(&IntExpr + 1)); - } + Expr *getIntExpr() { return hasIntExpr() ? getIntExprs()[0] : nullptr; }; }; class OpenACCNumWorkersClause : public OpenACCClauseWithSingleIntExpr { diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 1bbe76ff6bd2ac..a95424862e63f4 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12293,4 +12293,9 @@ def note_acc_int_expr_conversion : Note<"conversion to %select{integral|enumeration}0 type %1">; def err_acc_int_expr_multiple_conversions : Error<"multiple conversions from expression type %0 to an integral type">; +def err_acc_num_gangs_num_args + : Error<"%select{no|too many}0 integer expression arguments provided to " + "OpenACC 'num_gangs' " + "%select{|clause: '%1' directive expects maximum of %2, %3 were " + "provided}0">; } // end of sema component. diff --git a/clang/include/clang/Basic/OpenACCClauses.def b/clang/include/clang/Basic/OpenACCClauses.def index 520e068f4ffd40..dd5792e7ca8c39 100644 --- a/clang/include/clang/Basic/OpenACCClauses.def +++ b/clang/include/clang/Basic/OpenACCClauses.def @@ -18,6 +18,7 @@ VISIT_CLAUSE(Default) VISIT_CLAUSE(If) VISIT_CLAUSE(Self) +VISIT_CLAUSE(NumGangs) VISIT_CLAUSE(NumWorkers) VISIT_CLAUSE(VectorLength) diff --git a/clang/include/clang/Parse/Parser.h b/clang/include/clang/Parse/Parser.h index 72b2f958a5e622..d3bb04ff7a2c6d 100644 --- a/clang/include/clang/Parse/Parser.h +++ b/clang/include/clang/Parse/Parser.h @@ -3644,10 +3644,22 @@ class Parser : public CodeCompletionHandler { /// Parses the clause of the 'bind' argument, which can be a string literal or /// an ID expression. ExprResult ParseOpenACCBindClauseArgument(); + + /// A type to represent the state of parsing after an attempt to parse an + /// OpenACC int-expr. This is useful to determine whether an int-expr list can + /// continue parsing after a failed int-expr. + using OpenACCIntExprParseResult = + std::pair<ExprResult, OpenACCParseCanContinue>; /// Parses the clause kind of 'int-expr', which can be any integral /// expression. - ExprResult ParseOpenACCIntExpr(OpenACCDirectiveKind DK, OpenACCClauseKind CK, - SourceLocation Loc); + OpenACCIntExprParseResult ParseOpenACCIntExpr(OpenACCDirectiveKind DK, + OpenACCClauseKind CK, + SourceLocation Loc); + /// Parses the argument list for 'num_gangs', which allows up to 3 + /// 'int-expr's. + bool ParseOpenACCIntExprList(OpenACCDirectiveKind DK, OpenACCClauseKind CK, + SourceLocation Loc, + llvm::SmallVectorImpl<Expr *> &IntExprs); /// Parses the 'device-type-list', which is a list of identifiers. bool ParseOpenACCDeviceTypeList(); /// Parses the 'async-argument', which is an integral value with two diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h index 023722049732af..ea28617f79b81b 100644 --- a/clang/include/clang/Sema/SemaOpenACC.h +++ b/clang/include/clang/Sema/SemaOpenACC.h @@ -93,14 +93,16 @@ class SemaOpenACC : public SemaBase { } unsigned getNumIntExprs() const { - assert((ClauseKind == OpenACCClauseKind::NumWorkers || + assert((ClauseKind == OpenACCClauseKind::NumGangs || + ClauseKind == OpenACCClauseKind::NumWorkers || ClauseKind == OpenACCClauseKind::VectorLength) && "Parsed clause kind does not have a int exprs"); return std::get<IntExprDetails>(Details).IntExprs.size(); } ArrayRef<Expr *> getIntExprs() { - assert((ClauseKind == OpenACCClauseKind::NumWorkers || + assert((ClauseKind == OpenACCClauseKind::NumGangs || + ClauseKind == OpenACCClauseKind::NumWorkers || ClauseKind == OpenACCClauseKind::VectorLength) && "Parsed clause kind does not have a int exprs"); return std::get<IntExprDetails>(Details).IntExprs; @@ -134,11 +136,19 @@ class SemaOpenACC : public SemaBase { } void setIntExprDetails(ArrayRef<Expr *> IntExprs) { - assert((ClauseKind == OpenACCClauseKind::NumWorkers || + assert((ClauseKind == OpenACCClauseKind::NumGangs || + ClauseKind == OpenACCClauseKind::NumWorkers || ClauseKind == OpenACCClauseKind::VectorLength) && "Parsed clause kind does not have a int exprs"); Details = IntExprDetails{{IntExprs.begin(), IntExprs.end()}}; } + void setIntExprDetails(llvm::SmallVector<Expr *> &&IntExprs) { + assert((ClauseKind == OpenACCClauseKind::NumGangs || + ClauseKind == OpenACCClauseKind::NumWorkers || + ClauseKind == OpenACCClauseKind::VectorLength) && + "Parsed clause kind does not have a int exprs"); + Details = IntExprDetails{IntExprs}; + } }; SemaOpenACC(Sema &S); diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp index 75334223e073c3..6cd5b28802187d 100644 --- a/clang/lib/AST/OpenACCClause.cpp +++ b/clang/lib/AST/OpenACCClause.cpp @@ -124,6 +124,16 @@ OpenACCVectorLengthClause::Create(const ASTContext &C, SourceLocation BeginLoc, OpenACCVectorLengthClause(BeginLoc, LParenLoc, IntExpr, EndLoc); } +OpenACCNumGangsClause *OpenACCNumGangsClause::Create(const ASTContext &C, + SourceLocation BeginLoc, + SourceLocation LParenLoc, + ArrayRef<Expr *> IntExprs, + SourceLocation EndLoc) { + void *Mem = C.Allocate( + OpenACCNumGangsClause::totalSizeToAlloc<Expr *>(IntExprs.size())); + return new (Mem) OpenACCNumGangsClause(BeginLoc, LParenLoc, IntExprs, EndLoc); +} + //===----------------------------------------------------------------------===// // OpenACC clauses printing methods //===----------------------------------------------------------------------===// @@ -141,6 +151,12 @@ void OpenACCClausePrinter::VisitSelfClause(const OpenACCSelfClause &C) { OS << "(" << CondExpr << ")"; } +void OpenACCClausePrinter::VisitNumGangsClause(const OpenACCNumGangsClause &C) { + OS << "num_gangs("; + llvm::interleaveComma(C.getIntExprs(), OS); + OS << ")"; +} + void OpenACCClausePrinter::VisitNumWorkersClause( const OpenACCNumWorkersClause &C) { OS << "num_workers(" << C.getIntExpr() << ")"; diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp index 8138ae3a244a8b..c81724f84dd9ce 100644 --- a/clang/lib/AST/StmtProfile.cpp +++ b/clang/lib/AST/StmtProfile.cpp @@ -2497,6 +2497,12 @@ void OpenACCClauseProfiler::VisitSelfClause(const OpenACCSelfClause &Clause) { Profiler.VisitStmt(Clause.getConditionExpr()); } +void OpenACCClauseProfiler::VisitNumGangsClause( + const OpenACCNumGangsClause &Clause) { + for (auto *E : Clause.getIntExprs()) + Profiler.VisitStmt(E); +} + void OpenACCClauseProfiler::VisitNumWorkersClause( const OpenACCNumWorkersClause &Clause) { assert(Clause.hasIntExpr() && "num_workers clause requires a valid int expr"); diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp index e5a8b285715b30..8f0a9a9b0ed0bc 100644 --- a/clang/lib/AST/TextNodeDumper.cpp +++ b/clang/lib/AST/TextNodeDumper.cpp @@ -399,6 +399,7 @@ void TextNodeDumper::Visit(const OpenACCClause *C) { break; case OpenACCClauseKind::If: case OpenACCClauseKind::Self: + case OpenACCClauseKind::NumGangs: case OpenACCClauseKind::NumWorkers: case OpenACCClauseKind::VectorLength: // The condition expression will be printed as a part of the 'children', diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp index 757417f75c9636..8a18fca8064ee1 100644 --- a/clang/lib/Parse/ParseOpenACC.cpp +++ b/clang/lib/Parse/ParseOpenACC.cpp @@ -632,16 +632,54 @@ Parser::ParseOpenACCClauseList(OpenACCDirectiveKind DirKind) { return Clauses; } -ExprResult Parser::ParseOpenACCIntExpr(OpenACCDirectiveKind DK, - OpenACCClauseKind CK, - SourceLocation Loc) { - ExprResult ER = - getActions().CorrectDelayedTyposInExpr(ParseAssignmentExpression()); +Parser::OpenACCIntExprParseResult +Parser::ParseOpenACCIntExpr(OpenACCDirectiveKind DK, OpenACCClauseKind CK, + SourceLocation Loc) { + ExprResult ER = ParseAssignmentExpression(); + // If the actual parsing failed, we don't know the state of the parse, so + // don't try to continue. if (!ER.isUsable()) - return ER; + return {ER, OpenACCParseCanContinue::Cannot}; + + // Parsing can continue after the initial assignment expression parsing, so + // even if there was a typo, we can continue. + ER = getActions().CorrectDelayedTyposInExpr(ER); + if (!ER.isUsable()) + return {ER, OpenACCParseCanContinue::Can}; + + return {getActions().OpenACC().ActOnIntExpr(DK, CK, Loc, ER.get()), + OpenACCParseCanContinue::Can}; +} + +bool Parser::ParseOpenACCIntExprList(OpenACCDirectiveKind DK, + OpenACCClauseKind CK, SourceLocation Loc, + llvm::SmallVectorImpl<Expr *> &IntExprs) { + OpenACCIntExprParseResult CurResult = ParseOpenACCIntExpr(DK, CK, Loc); + + if (!CurResult.first.isUsable() && + CurResult.second == OpenACCParseCanContinue::Cannot) { + SkipUntil(tok::r_paren, tok::annot_pragma_openacc_end, + Parser::StopBeforeMatch); + return true; + } + + IntExprs.push_back(CurResult.first.get()); + + while (!getCurToken().isOneOf(tok::r_paren, tok::annot_pragma_openacc_end)) { + ExpectAndConsume(tok::comma); + + CurResult = ParseOpenACCIntExpr(DK, CK, Loc); - return getActions().OpenACC().ActOnIntExpr(DK, CK, Loc, ER.get()); + if (!CurResult.first.isUsable() && + CurResult.second == OpenACCParseCanContinue::Cannot) { + SkipUntil(tok::r_paren, tok::annot_pragma_openacc_end, + Parser::StopBeforeMatch); + return true; + } + IntExprs.push_back(CurResult.first.get()); + } + return false; } bool Parser::ParseOpenACCClauseVarList(OpenACCClauseKind Kind) { @@ -761,7 +799,7 @@ bool Parser::ParseOpenACCGangArg(SourceLocation GangLoc) { ConsumeToken(); return ParseOpenACCIntExpr(OpenACCDirectiveKind::Invalid, OpenACCClauseKind::Gang, GangLoc) - .isInvalid(); + .first.isInvalid(); } if (isOpenACCSpecialToken(OpenACCSpecialTokenKind::Num, getCurToken()) && @@ -773,7 +811,7 @@ bool Parser::ParseOpenACCGangArg(SourceLocation GangLoc) { // This is just the 'num' case where 'num' is optional. return ParseOpenACCIntExpr(OpenACCDirectiveKind::Invalid, OpenACCClauseKind::Gang, GangLoc) - .isInvalid(); + .first.isInvalid(); } bool Parser::ParseOpenACCGangArgList(SourceLocation GangLoc) { @@ -946,13 +984,25 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams( } break; } - case OpenACCClauseKind::NumGangs: + case OpenACCClauseKind::NumGangs: { + llvm::SmallVector<Expr *> IntExprs; + + if (ParseOpenACCIntExprList(OpenACCDirectiveKind::Invalid, + OpenACCClauseKind::NumGangs, ClauseLoc, + IntExprs)) { + Parens.skipToEnd(); + return OpenACCCanContinue(); + } + ParsedClause.setIntExprDetails(std::move(IntExprs)); + break; + } case OpenACCClauseKind::NumWorkers: case OpenACCClauseKind::DeviceNum: case OpenACCClauseKind::DefaultAsync: case OpenACCClauseKind::VectorLength: { ExprResult IntExpr = ParseOpenACCIntExpr(OpenACCDirectiveKind::Invalid, - ClauseKind, ClauseLoc); + ClauseKind, ClauseLoc) + .first; if (IntExpr.isInvalid()) { Parens.skipToEnd(); return OpenACCCanContinue(); @@ -1017,7 +1067,8 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams( : OpenACCSpecialTokenKind::Num, ClauseKind); ExprResult IntExpr = ParseOpenACCIntExpr(OpenACCDirectiveKind::Invalid, - ClauseKind, ClauseLoc); + ClauseKind, ClauseLoc) + .first; if (IntExpr.isInvalid()) { Parens.skipToEnd(); return OpenACCCanContinue(); @@ -1081,11 +1132,13 @@ bool Parser::ParseOpenACCWaitArgument(SourceLocation Loc, bool IsDirective) { // Consume colon. ConsumeToken(); - ExprResult IntExpr = ParseOpenACCIntExpr( - IsDirective ? OpenACCDirectiveKind::Wait - : OpenACCDirectiveKind::Invalid, - IsDirective ? OpenACCClauseKind::Invalid : OpenACCClauseKind::Wait, - Loc); + ExprResult IntExpr = + ParseOpenACCIntExpr(IsDirective ? OpenACCDirectiveKind::Wait + : OpenACCDirectiveKind::Invalid, + IsDirective ? OpenACCClauseKind::Invalid + : OpenACCClauseKind::Wait, + Loc) + .first; if (IntExpr.isInvalid()) return true; diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp index 190739fa02e932..ba69e71e30a181 100644 --- a/clang/lib/Sema/SemaOpenACC.cpp +++ b/clang/lib/Sema/SemaOpenACC.cpp @@ -91,6 +91,7 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind, default: return false; } + case OpenACCClauseKind::NumGangs: case OpenACCClauseKind::NumWorkers: case OpenACCClauseKind::VectorLength: switch (DirectiveKind) { @@ -230,6 +231,40 @@ SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses, getASTContext(), Clause.getBeginLoc(), Clause.getLParenLoc(), Clause.getConditionExpr(), Clause.getEndLoc()); } + case OpenACCClauseKind::NumGangs: { + // Restrictions only properly implemented on 'compute' constructs, and + // 'compute' constructs are the only construct that can do anything with + // this yet, so skip/treat as unimplemented in this case. + if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind())) + break; + + // 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(*this, ExistingClauses, Clause)) + return nullptr; + + if (Clause.getIntExprs().empty()) + Diag(Clause.getBeginLoc(), diag::err_acc_num_gangs_num_args) + << /*NoArgs=*/0; + + unsigned MaxArgs = + (Clause.getDirectiveKind() == OpenACCDirectiveKind::Parallel || + Clause.getDirectiveKind() == OpenACCDirectiveKind::ParallelLoop) + ? 3 + : 1; + if (Clause.getIntExprs().size() > MaxArgs) + Diag(Clause.getBeginLoc(), diag::err_acc_num_gangs_num_args) + << /*NoArgs=*/1 << Clause.getDirectiveKind() << MaxArgs + << Clause.getIntExprs().size(); + + // Create the AST node for the clause even if the number of expressions is + // incorrect. + return OpenACCNumGangsClause::Create( + getASTContext(), Clause.getBeginLoc(), Clause.getLParenLoc(), + Clause.getIntExprs(), Clause.getEndLoc()); + break; + } case OpenACCClauseKind::NumWorkers: { // Restrictions only properly implemented on 'compute' constructs, and // 'compute' constructs are the only construct that can do anything with diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h index 13f7e9b3fbe365..9404be5a46f3f7 100644 --- a/clang/lib/Sema/TreeTransform.h +++ b/clang/lib/Sema/TreeTransform.h @@ -11162,6 +11162,32 @@ void OpenACCClauseTransform<Derived>::VisitSelfClause( ParsedClause.getEndLoc()); } +template <typename Derived> +void OpenACCClauseTransform<Derived>::VisitNumGangsClause( + const OpenACCNumGangsClause &C) { + llvm::SmallVector<Expr *> InstantiatedIntExprs; + + for (Expr *CurIntExpr : C.getIntExprs()) { + ExprResult Res = Self.TransformExpr(CurIntExpr); + + if (!Res.isUsable()) + return; + + Res = Self.getSema().OpenACC().ActOnIntExpr(OpenACCDirectiveKind::Invalid, + C.getClauseKind(), + C.getBeginLoc(), Res.get()); + if (!Res.isUsable()) + return; + + InstantiatedIntExprs.push_back(Res.get()); + } + + ParsedClause.setIntExprDetails(InstantiatedIntExprs); + NewClause = OpenACCNumGangsClause::Create( + Self.getSema().getASTContext(), ParsedClause.getBeginLoc(), + ParsedClause.getLParenLoc(), ParsedClause.getIntExprs(), + ParsedClause.getEndLoc()); +} template <typename Derived> void OpenACCClauseTransform<Derived>::VisitNumWorkersClause( const OpenACCNumWorkersClause &C) { diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp index 44e23919ea18e0..d64925676df7b1 100644 --- a/clang/lib/Serialization/ASTReader.cpp +++ b/clang/lib/Serialization/ASTReader.cpp @@ -11786,6 +11786,15 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() { return OpenACCSelfClause::Create(getContext(), BeginLoc, LParenLoc, CondExpr, EndLoc); } + case OpenACCClauseKind::NumGangs: { + SourceLocation LParenLoc = readSourceLocation(); + unsigned NumClauses = readInt(); + llvm::SmallVector<Expr *> IntExprs; + for (unsigned I = 0; I < NumClauses; ++I) + IntExprs.push_back(readSubExpr()); + return OpenACCNumGangsClause::Create(getContext(), BeginLoc, LParenLoc, + IntExprs, EndLoc); + } case OpenACCClauseKind::NumWorkers: { SourceLocation LParenLoc = readSourceLocation(); Expr *IntExpr = readSubExpr(); @@ -11826,7 +11835,6 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() { case OpenACCClauseKind::Reduction: case OpenACCClauseKind::Collapse: case OpenACCClauseKind::Bind: - case OpenACCClauseKind::NumGangs: case OpenACCClauseKind::DeviceNum: case OpenACCClauseKind::DefaultAsync: case OpenACCClauseKind::DeviceType: diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp index 8a4b36207c4734..018b854652a46e 100644 --- a/clang/lib/Serialization/ASTWriter.cpp +++ b/clang/lib/Serialization/ASTWriter.cpp @@ -7657,6 +7657,14 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) { AddStmt(const_cast<Expr *>(SC->getConditionExpr())); return; } + case OpenACCClauseKind::NumGangs: { + const auto *NGC = cast<OpenACCNumGangsClause>(C); + writeSourceLocation(NGC->getLParenLoc()); + writeUInt32(NGC->getIntExprs().size()); + for (Expr *E : NGC->getIntExprs()) + AddStmt(E); + return; + } case OpenACCClauseKind::NumWorkers: { const auto *NWC = cast<OpenACCNumWorkersClause>(C); writeSourceLocation(NWC->getLParenLoc()); @@ -7697,7 +7705,6 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) { case OpenACCClauseKind::Reduction: case OpenACCClauseKind::Collapse: case OpenACCClauseKind::Bind: - case OpenACCClauseKind::NumGangs: case OpenACCClauseKind::DeviceNum: case OpenACCClauseKind::DefaultAsync: case OpenACCClauseKind::DeviceType: diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c index ddf40f71701eda..799f22b8c120e5 100644 --- a/clang/test/ParserOpenACC/parse-clauses.c +++ b/clang/test/ParserOpenACC/parse-clauses.c @@ -911,16 +911,12 @@ void IntExprParsing() { #pragma acc parallel num_gangs(invalid) {} - // expected-error@+2{{expected ')'}} - // expected-note@+1{{to match this '('}} #pragma acc parallel num_gangs(5, 4) {} - // expected-warning@+1{{OpenACC clause 'num_gangs' not yet implemented, clause ignored}} #pragma acc parallel num_gangs(5) {} - // expected-warning@+1{{OpenACC clause 'num_gangs' not yet implemented, clause ignored}} #pragma acc parallel num_gangs(returns_int()) {} diff --git a/clang/test/SemaOpenACC/compute-construct-intexpr-clause-ast.cpp b/clang/test/SemaOpenACC/compute-construct-intexpr-clause-ast.cpp index 889025c26818d8..5a4c9f05ee089e 100644 --- a/clang/test/SemaOpenACC/compute-construct-intexpr-clause-ast.cpp +++ b/clang/test/SemaOpenACC/compute-construct-intexpr-clause-ast.cpp @@ -88,6 +88,34 @@ void NormalUses() { // CHECK-NEXT: WhileStmt // CHECK-NEXT: CXXBoolLiteralExpr // CHECK-NEXT: CompoundStmt + +#pragma acc parallel num_gangs(some_int(), some_long(), some_short()) + while(true){} + // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel + // CHECK-NEXT: num_gangs clause + // CHECK-NEXT: CallExpr{{.*}}'int' + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int (*)()' <FunctionToPointerDecay> + // CHECK-NEXT: DeclRefExpr{{.*}}'int ()' lvalue Function{{.*}} 'some_int' 'int ()' + // CHECK-NEXT: CallExpr{{.*}}'long' + // CHECK-NEXT: ImplicitCastExpr{{.*}}'long (*)()' <FunctionToPointerDecay> + // CHECK-NEXT: DeclRefExpr{{.*}}'long ()' lvalue Function{{.*}} 'some_long' 'long ()' + // CHECK-NEXT: CallExpr{{.*}}'short' + // CHECK-NEXT: ImplicitCastExpr{{.*}}'short (*)()' <FunctionToPointerDecay> + // CHECK-NEXT: DeclRefExpr{{.*}}'short ()' lvalue Function{{.*}} 'some_short' 'short ()' + // CHECK-NEXT: WhileStmt + // CHECK-NEXT: CXXBoolLiteralExpr + // CHECK-NEXT: CompoundStmt + +#pragma acc kernels num_gangs(some_int()) + while(true){} + // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels + // CHECK-NEXT: num_gangs clause + // CHECK-NEXT: CallExpr{{.*}}'int' + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int (*)()' <FunctionToPointerDecay> + // CHECK-NEXT: DeclRefExpr{{.*}}'int ()' lvalue Function{{.*}} 'some_int' 'int ()' + // CHECK-NEXT: WhileStmt + // CHECK-NEXT: CXXBoolLiteralExpr + // CHECK-NEXT: CompoundStmt } template<typename T, typename U> @@ -187,6 +215,31 @@ void TemplUses(T t, U u) { // CHECK-NEXT: CXXBoolLiteralExpr // CHECK-NEXT: CompoundStmt +#pragma acc kernels num_gangs(u) + while(true){} + // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels + // CHECK-NEXT: num_gangs clause + // CHECK-NEXT: DeclRefExpr{{.*}} 'U' lvalue ParmVar{{.*}} 'u' 'U' + // CHECK-NEXT: WhileStmt + // CHECK-NEXT: CXXBoolLiteralExpr + // CHECK-NEXT: CompoundStmt + +#pragma acc parallel num_gangs(u, U::value) + while(true){} + // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel + // CHECK-NEXT: num_gangs clause + // CHECK-NEXT: DeclRefExpr{{.*}} 'U' lvalue ParmVar{{.*}} 'u' 'U' + // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue + // CHECK-NEXT: NestedNameSpecifier TypeSpec 'U' + // CHECK-NEXT: WhileStmt + // CHECK-NEXT: CXXBoolLiteralExpr + // CHECK-NEXT: CompoundStmt + + + // CHECK-NEXT: DeclStmt + // CHECK-NEXT: VarDecl{{.*}}EndMarker + int EndMarker; + // Check the instantiated versions of the above. // CHECK-NEXT: FunctionDecl{{.*}} used TemplUses 'void (CorrectConvert, HasInt)' implicit_instantiation // CHECK-NEXT: TemplateArgument type 'CorrectConvert' @@ -288,6 +341,32 @@ void TemplUses(T t, U u) { // CHECK-NEXT: WhileStmt // CHECK-NEXT: CXXBoolLiteralExpr // CHECK-NEXT: CompoundStmt + + // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels + // CHECK-NEXT: num_gangs clause + // CHECK-NEXT: ImplicitCastExpr{{.*}} 'char' <UserDefinedConversion> + // CHECK-NEXT: CXXMemberCallExpr{{.*}}'char' + // CHECK-NEXT: MemberExpr{{.*}} '<bound member function type>' .operator char + // CHECK-NEXT: DeclRefExpr{{.*}} 'HasInt' lvalue ParmVar + // CHECK-NEXT: WhileStmt + // CHECK-NEXT: CXXBoolLiteralExpr + // CHECK-NEXT: CompoundStmt + + // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel + // CHECK-NEXT: num_gangs clause + // CHECK-NEXT: ImplicitCastExpr{{.*}} 'char' <UserDefinedConversion> + // CHECK-NEXT: CXXMemberCallExpr{{.*}}'char' + // CHECK-NEXT: MemberExpr{{.*}} '<bound member function type>' .operator char + // CHECK-NEXT: DeclRefExpr{{.*}} 'HasInt' lvalue ParmVar + // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue> + // CHECK-NEXT: DeclRefExpr{{.*}} 'const int' lvalue Var{{.*}} 'value' 'const int' + // CHECK-NEXT: NestedNameSpecifier TypeSpec 'HasInt' + // CHECK-NEXT: WhileStmt + // CHECK-NEXT: CXXBoolLiteralExpr + // CHECK-NEXT: CompoundStmt + + // CHECK-NEXT: DeclStmt + // CHECK-NEXT: VarDecl{{.*}}EndMarker } struct HasInt { diff --git a/clang/test/SemaOpenACC/compute-construct-num_gangs-clause.c b/clang/test/SemaOpenACC/compute-construct-num_gangs-clause.c new file mode 100644 index 00000000000000..cdc6847b47f948 --- /dev/null +++ b/clang/test/SemaOpenACC/compute-construct-num_gangs-clause.c @@ -0,0 +1,54 @@ +// RUN: %clang_cc1 %s -fopenacc -verify + +short getS(); +void Test() { +#pragma acc kernels num_gangs(1) + while(1); + + // expected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}} +#pragma acc serial num_gangs(1) + while(1); + +#pragma acc parallel num_gangs(1) + while(1); + + // expected-error@+2{{OpenACC 'num_gangs' clause cannot appear more than once on a 'kernels' directive}} + // expected-note@+1{{previous clause is here}} +#pragma acc kernels num_gangs(1) num_gangs(2) + while(1); + + // expected-error@+2{{OpenACC 'num_gangs' clause cannot appear more than once on a 'parallel' directive}} + // expected-note@+1{{previous clause is here}} +#pragma acc parallel num_gangs(1) 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); + + // expected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}} +#pragma acc serial num_gangs(1, getS()) + while(1); +#pragma acc parallel num_gangs(1, getS()) + while(1); + + struct NotConvertible{} NC; + // expected-error@+1{{OpenACC clause 'num_gangs' requires expression of integer type ('struct NotConvertible' invalid)}} +#pragma acc parallel num_gangs(NC) + while(1); + + // expected-error@+1{{OpenACC clause 'num_gangs' requires expression of integer type ('struct NotConvertible' invalid)}} +#pragma acc parallel num_gangs(1, NC) + while(1); + + // expected-error@+1{{OpenACC clause 'num_gangs' requires expression of integer type ('struct NotConvertible' invalid)}} +#pragma acc parallel num_gangs(NC, 1) + while(1); + +#pragma acc parallel num_gangs(getS(), 1, getS()) + while(1); + + // expected-error@+1{{too many integer expression arguments provided to OpenACC 'num_gangs' clause: 'parallel' directive expects maximum of 3, 4 were provided}} +#pragma acc parallel num_gangs(getS(), 1, getS(), 1) + while(1); +} diff --git a/clang/test/SemaOpenACC/compute-construct-num_gangs-clause.cpp b/clang/test/SemaOpenACC/compute-construct-num_gangs-clause.cpp new file mode 100644 index 00000000000000..ec3df87a065572 --- /dev/null +++ b/clang/test/SemaOpenACC/compute-construct-num_gangs-clause.cpp @@ -0,0 +1,160 @@ +// RUN: %clang_cc1 %s -fopenacc -verify + +struct NotConvertible{} NC; +struct Incomplete *SomeIncomplete; // #INCOMPLETE +enum E{} SomeE; +enum class E2{} SomeE2; + +struct CorrectConvert { + operator int(); +} Convert; + +struct ExplicitConvertOnly { + explicit operator int() const; // #EXPL_CONV +} Explicit; + +struct AmbiguousConvert{ + operator int(); // #AMBIG_INT + operator short(); // #AMBIG_SHORT + operator float(); +} Ambiguous; + +short some_short(); +int some_int(); +long some_long(); + +void Test() { +#pragma acc kernels num_gangs(1) + while(1); + + // expected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}} +#pragma acc serial num_gangs(1) + while(1); + +#pragma acc parallel num_gangs(1) + while(1); + + // expected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}} +#pragma acc serial num_gangs(some_short(), some_int(), some_long()) + while(1); + +#pragma acc parallel num_gangs(some_short(), some_int(), some_long()) + while(1); + + // expected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}} +#pragma acc serial num_gangs(some_short(), some_int(), some_long(), SomeE) + while(1); + + // expected-error@+1{{too many integer expression arguments provided to OpenACC 'num_gangs' clause: 'parallel' directive expects maximum of 3, 4 were provided}} +#pragma acc parallel num_gangs(some_short(), some_int(), some_long(), SomeE) + 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, 2) + while(1); + + // expected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}} +#pragma acc serial num_gangs(1, 2) + while(1); + +#pragma acc parallel num_gangs(1, 2) + while(1); + + // 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'}} +#pragma acc parallel num_gangs(Ambiguous) + while(1); + + // expected-error@+1{{OpenACC clause 'num_gangs' requires expression of integer type ('struct NotConvertible' invalid)}} +#pragma acc parallel num_gangs(NC, SomeE) + while(1); + + // expected-error@+1{{OpenACC clause 'num_gangs' requires expression of integer type ('struct NotConvertible' invalid)}} +#pragma acc parallel num_gangs(SomeE, NC) + while(1); + + // expected-error@+3{{OpenACC integer expression type 'struct ExplicitConvertOnly' requires explicit conversion to 'int'}} + // expected-note@#EXPL_CONV{{conversion to integral type 'int'}} + // expected-error@+1{{OpenACC clause 'num_gangs' requires expression of integer type ('struct NotConvertible' invalid)}} +#pragma acc parallel num_gangs(Explicit, NC) + while(1); + + // expected-error@+4{{OpenACC integer expression type 'struct ExplicitConvertOnly' requires explicit conversion to 'int'}} + // expected-note@#EXPL_CONV{{conversion to integral type 'int'}} + // expected-error@+2{{OpenACC clause 'num_gangs' requires expression of integer type ('struct NotConvertible' invalid)}} + // expected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}} +#pragma acc serial num_gangs(Explicit, NC) + while(1); + + // expected-error@+6{{OpenACC integer expression type 'struct ExplicitConvertOnly' requires explicit conversion to 'int'}} + // expected-note@#EXPL_CONV{{conversion to integral type 'int'}} + // expected-error@+4{{OpenACC clause 'num_gangs' requires expression of integer type ('struct NotConvertible' invalid)}} + // 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'}} +#pragma acc parallel num_gangs(Explicit, NC, Ambiguous) + while(1); + + // expected-error@+7{{OpenACC integer expression type 'struct ExplicitConvertOnly' requires explicit conversion to 'int'}} + // expected-note@#EXPL_CONV{{conversion to integral type 'int'}} + // expected-error@+5{{OpenACC clause 'num_gangs' requires expression of integer type ('struct NotConvertible' invalid)}} + // 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 'num_gangs' clause is not valid on 'serial' directive}} +#pragma acc serial num_gangs(Explicit, NC, Ambiguous) + while(1); + // TODO +} + +struct HasInt { + using IntTy = int; + using ShortTy = short; + static constexpr int value = 1; + static constexpr AmbiguousConvert ACValue; + static constexpr ExplicitConvertOnly EXValue; + + operator char(); +}; + +template <typename T> +void TestInst() { + // expected-error@+2{{no member named 'Invalid' in 'HasInt'}} + // expected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}} +#pragma acc serial num_gangs(HasInt::Invalid) + while(1); + + // expected-error@+2{{no member named 'Invalid' in 'HasInt'}} + // expected-note@#INST{{in instantiation of function template specialization}} +#pragma acc parallel num_gangs(T::Invalid) + while(1); + + // expected-error@+1{{no member named 'Invalid' in 'HasInt'}} +#pragma acc parallel num_gangs(1, HasInt::Invalid) + while(1); + + // expected-error@+1{{no member named 'Invalid' in 'HasInt'}} +#pragma acc parallel num_gangs(T::Invalid, 1) + while(1); + + // expected-error@+2{{no member named 'Invalid' in 'HasInt'}} + // expected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}} +#pragma acc serial num_gangs(1, HasInt::Invalid) + while(1); + + // expected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}} +#pragma acc serial num_gangs(T::Invalid, 1) + while(1); + +#pragma acc parallel num_gangs(T::value, typename T::IntTy{}) + while(1); + + // expected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}} +#pragma acc serial num_gangs(T::value, typename T::IntTy{}) + while(1); +} + +void Inst() { + TestInst<HasInt>(); // #INST +} diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp index cbc1d85bb33dfc..74163f30e19b1d 100644 --- a/clang/tools/libclang/CIndex.cpp +++ b/clang/tools/libclang/CIndex.cpp @@ -2803,6 +2803,10 @@ void OpenACCClauseEnqueue::VisitVectorLengthClause( const OpenACCVectorLengthClause &C) { Visitor.AddStmt(C.getIntExpr()); } +void OpenACCClauseEnqueue::VisitNumGangsClause(const OpenACCNumGangsClause &C) { + for (Expr *IE : C.getIntExprs()) + Visitor.AddStmt(IE); +} } // 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
