Author: Erich Keane Date: 2024-10-03T08:34:43-07:00 New Revision: d412cea8c4f26f451aee46641e384e8df62a5904
URL: https://github.com/llvm/llvm-project/commit/d412cea8c4f26f451aee46641e384e8df62a5904 DIFF: https://github.com/llvm/llvm-project/commit/d412cea8c4f26f451aee46641e384e8df62a5904.diff LOG: [OpenACC] Implement 'tile' attribute AST (#110999) The 'tile' clause shares quite a bit of the rules with 'collapse', so a followup patch will add those tests/behaviors. This patch deals with adding the AST node. The 'tile' clause takes a series of integer constant expressions, or *. The asterisk is now represented by a new OpenACCAsteriskSizeExpr node, else this clause is very similar to others. Added: clang/test/SemaOpenACC/loop-construct-tile-ast.cpp clang/test/SemaOpenACC/loop-construct-tile-clause.cpp Modified: clang/include/clang/AST/ComputeDependence.h clang/include/clang/AST/Expr.h clang/include/clang/AST/JSONNodeDumper.h clang/include/clang/AST/OpenACCClause.h clang/include/clang/AST/RecursiveASTVisitor.h clang/include/clang/AST/TextNodeDumper.h clang/include/clang/Basic/DiagnosticSemaKinds.td clang/include/clang/Basic/OpenACCClauses.def clang/include/clang/Basic/StmtNodes.td clang/include/clang/Parse/Parser.h clang/include/clang/Sema/SemaOpenACC.h clang/include/clang/Serialization/ASTBitCodes.h clang/lib/AST/ComputeDependence.cpp clang/lib/AST/Expr.cpp clang/lib/AST/ExprClassification.cpp clang/lib/AST/ExprConstant.cpp clang/lib/AST/ItaniumMangle.cpp clang/lib/AST/JSONNodeDumper.cpp clang/lib/AST/OpenACCClause.cpp clang/lib/AST/StmtPrinter.cpp clang/lib/AST/StmtProfile.cpp clang/lib/AST/TextNodeDumper.cpp clang/lib/CodeGen/CGExprScalar.cpp clang/lib/Parse/ParseOpenACC.cpp clang/lib/Sema/SemaExceptionSpec.cpp clang/lib/Sema/SemaOpenACC.cpp clang/lib/Sema/TreeTransform.h clang/lib/Serialization/ASTReader.cpp clang/lib/Serialization/ASTReaderStmt.cpp clang/lib/Serialization/ASTWriter.cpp clang/lib/Serialization/ASTWriterStmt.cpp clang/lib/StaticAnalyzer/Core/ExprEngine.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 clang/tools/libclang/CXCursor.cpp Removed: ################################################################################ diff --git a/clang/include/clang/AST/ComputeDependence.h b/clang/include/clang/AST/ComputeDependence.h index 6d3a51c379f9df..1a8507cfbf9872 100644 --- a/clang/include/clang/AST/ComputeDependence.h +++ b/clang/include/clang/AST/ComputeDependence.h @@ -107,6 +107,7 @@ class ObjCSubscriptRefExpr; class ObjCIsaExpr; class ObjCIndirectCopyRestoreExpr; class ObjCMessageExpr; +class OpenACCAsteriskSizeExpr; // The following functions are called from constructors of `Expr`, so they // should not access anything beyond basic @@ -203,6 +204,7 @@ ExprDependence computeDependence(ObjCSubscriptRefExpr *E); ExprDependence computeDependence(ObjCIsaExpr *E); ExprDependence computeDependence(ObjCIndirectCopyRestoreExpr *E); ExprDependence computeDependence(ObjCMessageExpr *E); +ExprDependence computeDependence(OpenACCAsteriskSizeExpr *E); } // namespace clang #endif diff --git a/clang/include/clang/AST/Expr.h b/clang/include/clang/AST/Expr.h index 66c746cc25040f..57353855c51e7c 100644 --- a/clang/include/clang/AST/Expr.h +++ b/clang/include/clang/AST/Expr.h @@ -2072,6 +2072,41 @@ class PredefinedExpr final } }; +/// This expression type represents an asterisk in an OpenACC Size-Expr, used in +/// the 'tile' and 'gang' clauses. It is of 'int' type, but should not be +/// evaluated. +class OpenACCAsteriskSizeExpr final : public Expr { + friend class ASTStmtReader; + SourceLocation AsteriskLoc; + + OpenACCAsteriskSizeExpr(SourceLocation AsteriskLoc, QualType IntTy) + : Expr(OpenACCAsteriskSizeExprClass, IntTy, VK_PRValue, OK_Ordinary), + AsteriskLoc(AsteriskLoc) {} + + void setAsteriskLocation(SourceLocation Loc) { AsteriskLoc = Loc; } + +public: + static OpenACCAsteriskSizeExpr *Create(const ASTContext &C, + SourceLocation Loc); + static OpenACCAsteriskSizeExpr *CreateEmpty(const ASTContext &C); + + SourceLocation getBeginLoc() const { return AsteriskLoc; } + SourceLocation getEndLoc() const { return AsteriskLoc; } + SourceLocation getLocation() const { return AsteriskLoc; } + + static bool classof(const Stmt *T) { + return T->getStmtClass() == OpenACCAsteriskSizeExprClass; + } + // Iterators + 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()); + } +}; + // This represents a use of the __builtin_sycl_unique_stable_name, which takes a // type-id, and at CodeGen time emits a unique string representation of the // type in a way that permits us to properly encode information about the SYCL diff --git a/clang/include/clang/AST/JSONNodeDumper.h b/clang/include/clang/AST/JSONNodeDumper.h index 55bd583e304e8b..9422c8fceccfbd 100644 --- a/clang/include/clang/AST/JSONNodeDumper.h +++ b/clang/include/clang/AST/JSONNodeDumper.h @@ -283,6 +283,7 @@ class JSONNodeDumper void VisitDeclRefExpr(const DeclRefExpr *DRE); void VisitSYCLUniqueStableNameExpr(const SYCLUniqueStableNameExpr *E); + void VisitOpenACCAsteriskSizeExpr(const OpenACCAsteriskSizeExpr *E); void VisitPredefinedExpr(const PredefinedExpr *PE); void VisitUnaryOperator(const UnaryOperator *UO); void VisitBinaryOperator(const BinaryOperator *BO); diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h index 90f5b7fc9ab6f4..e4f2e07222a338 100644 --- a/clang/include/clang/AST/OpenACCClause.h +++ b/clang/include/clang/AST/OpenACCClause.h @@ -481,6 +481,35 @@ class OpenACCNumGangsClause final } }; +class OpenACCTileClause final + : public OpenACCClauseWithExprs, + public llvm::TrailingObjects<OpenACCTileClause, Expr *> { + OpenACCTileClause(SourceLocation BeginLoc, SourceLocation LParenLoc, + ArrayRef<Expr *> SizeExprs, SourceLocation EndLoc) + : OpenACCClauseWithExprs(OpenACCClauseKind::Tile, BeginLoc, LParenLoc, + EndLoc) { + std::uninitialized_copy(SizeExprs.begin(), SizeExprs.end(), + getTrailingObjects<Expr *>()); + setExprs(MutableArrayRef(getTrailingObjects<Expr *>(), SizeExprs.size())); + } + +public: + static bool classof(const OpenACCClause *C) { + return C->getClauseKind() == OpenACCClauseKind::Tile; + } + static OpenACCTileClause *Create(const ASTContext &C, SourceLocation BeginLoc, + SourceLocation LParenLoc, + ArrayRef<Expr *> SizeExprs, + SourceLocation EndLoc); + llvm::ArrayRef<Expr *> getSizeExprs() { + return OpenACCClauseWithExprs::getExprs(); + } + + llvm::ArrayRef<Expr *> getSizeExprs() const { + return OpenACCClauseWithExprs::getExprs(); + } +}; + /// Represents one of a handful of clauses that have a single integer /// expression. class OpenACCClauseWithSingleIntExpr : public OpenACCClauseWithExprs { diff --git a/clang/include/clang/AST/RecursiveASTVisitor.h b/clang/include/clang/AST/RecursiveASTVisitor.h index cd9947f7ab9805..cbbba9e88b7f5b 100644 --- a/clang/include/clang/AST/RecursiveASTVisitor.h +++ b/clang/include/clang/AST/RecursiveASTVisitor.h @@ -2867,6 +2867,7 @@ DEF_TRAVERSE_STMT(ParenListExpr, {}) DEF_TRAVERSE_STMT(SYCLUniqueStableNameExpr, { TRY_TO(TraverseTypeLoc(S->getTypeSourceInfo()->getTypeLoc())); }) +DEF_TRAVERSE_STMT(OpenACCAsteriskSizeExpr, {}) DEF_TRAVERSE_STMT(PredefinedExpr, {}) DEF_TRAVERSE_STMT(ShuffleVectorExpr, {}) DEF_TRAVERSE_STMT(ConvertVectorExpr, {}) diff --git a/clang/include/clang/AST/TextNodeDumper.h b/clang/include/clang/AST/TextNodeDumper.h index 57100e7ede171c..9c320c8ae3e54c 100644 --- a/clang/include/clang/AST/TextNodeDumper.h +++ b/clang/include/clang/AST/TextNodeDumper.h @@ -410,6 +410,7 @@ class TextNodeDumper void VisitHLSLOutArgExpr(const HLSLOutArgExpr *E); void VisitOpenACCConstructStmt(const OpenACCConstructStmt *S); void VisitOpenACCLoopConstruct(const OpenACCLoopConstruct *S); + void VisitOpenACCAsteriskSizeExpr(const OpenACCAsteriskSizeExpr *S); void VisitEmbedExpr(const EmbedExpr *S); void VisitAtomicExpr(const AtomicExpr *AE); }; diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index dc84110ef78211..552120f6b7c9d3 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12666,6 +12666,10 @@ def err_acc_loop_spec_conflict def err_acc_collapse_loop_count : Error<"OpenACC 'collapse' clause loop count must be a %select{constant " "expression|positive integer value, evaluated to %1}0">; +def err_acc_size_expr_value + : Error< + "OpenACC 'tile' clause size expression must be %select{an asterisk " + "or a constant expression|positive integer value, evaluated to %1}0">; def err_acc_invalid_in_collapse_loop : Error<"%select{OpenACC '%1' construct|while loop|do loop}0 cannot appear " "in intervening code of a 'loop' with a 'collapse' clause">; diff --git a/clang/include/clang/Basic/OpenACCClauses.def b/clang/include/clang/Basic/OpenACCClauses.def index 19cdfe7672133b..a380e5ae69c418 100644 --- a/clang/include/clang/Basic/OpenACCClauses.def +++ b/clang/include/clang/Basic/OpenACCClauses.def @@ -52,6 +52,7 @@ VISIT_CLAUSE(Private) VISIT_CLAUSE(Reduction) VISIT_CLAUSE(Self) VISIT_CLAUSE(Seq) +VISIT_CLAUSE(Tile) VISIT_CLAUSE(VectorLength) VISIT_CLAUSE(Wait) diff --git a/clang/include/clang/Basic/StmtNodes.td b/clang/include/clang/Basic/StmtNodes.td index 30f2c8f1dbfde8..dc82a4b6d1f777 100644 --- a/clang/include/clang/Basic/StmtNodes.td +++ b/clang/include/clang/Basic/StmtNodes.td @@ -308,5 +308,8 @@ def OpenACCAssociatedStmtConstruct def OpenACCComputeConstruct : StmtNode<OpenACCAssociatedStmtConstruct>; def OpenACCLoopConstruct : StmtNode<OpenACCAssociatedStmtConstruct>; +// OpenACC Additional Expressions. +def OpenACCAsteriskSizeExpr : StmtNode<Expr>; + // HLSL Constructs. def HLSLOutArgExpr : StmtNode<Expr>; diff --git a/clang/include/clang/Parse/Parser.h b/clang/include/clang/Parse/Parser.h index eb8a851da7e04e..93e49d395388a6 100644 --- a/clang/include/clang/Parse/Parser.h +++ b/clang/include/clang/Parse/Parser.h @@ -3786,10 +3786,13 @@ class Parser : public CodeCompletionHandler { OpenACCIntExprParseResult ParseOpenACCAsyncArgument(OpenACCDirectiveKind DK, OpenACCClauseKind CK, SourceLocation Loc); + /// Parses the 'size-expr', which is an integral value, or an asterisk. - bool ParseOpenACCSizeExpr(); + /// Asterisk is represented by a OpenACCAsteriskSizeExpr + ExprResult ParseOpenACCSizeExpr(OpenACCClauseKind CK); /// Parses a comma delimited list of 'size-expr's. - bool ParseOpenACCSizeExprList(); + 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. diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h index 26564835fa1af6..d25c52ec3be43a 100644 --- a/clang/include/clang/Sema/SemaOpenACC.h +++ b/clang/include/clang/Sema/SemaOpenACC.h @@ -179,6 +179,7 @@ class SemaOpenACC : public SemaBase { assert((ClauseKind == OpenACCClauseKind::NumGangs || ClauseKind == OpenACCClauseKind::NumWorkers || ClauseKind == OpenACCClauseKind::Async || + ClauseKind == OpenACCClauseKind::Tile || ClauseKind == OpenACCClauseKind::VectorLength) && "Parsed clause kind does not have a int exprs"); @@ -224,6 +225,7 @@ class SemaOpenACC : public SemaBase { assert((ClauseKind == OpenACCClauseKind::NumGangs || ClauseKind == OpenACCClauseKind::NumWorkers || ClauseKind == OpenACCClauseKind::Async || + ClauseKind == OpenACCClauseKind::Tile || ClauseKind == OpenACCClauseKind::VectorLength) && "Parsed clause kind does not have a int exprs"); @@ -335,6 +337,7 @@ class SemaOpenACC : public SemaBase { assert((ClauseKind == OpenACCClauseKind::NumGangs || ClauseKind == OpenACCClauseKind::NumWorkers || ClauseKind == OpenACCClauseKind::Async || + ClauseKind == OpenACCClauseKind::Tile || ClauseKind == OpenACCClauseKind::VectorLength) && "Parsed clause kind does not have a int exprs"); Details = IntExprDetails{{IntExprs.begin(), IntExprs.end()}}; @@ -343,6 +346,7 @@ class SemaOpenACC : public SemaBase { assert((ClauseKind == OpenACCClauseKind::NumGangs || ClauseKind == OpenACCClauseKind::NumWorkers || ClauseKind == OpenACCClauseKind::Async || + ClauseKind == OpenACCClauseKind::Tile || ClauseKind == OpenACCClauseKind::VectorLength) && "Parsed clause kind does not have a int exprs"); Details = IntExprDetails{std::move(IntExprs)}; @@ -522,6 +526,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. + ExprResult CheckTileSizeExpr(Expr *SizeExpr); + + ExprResult BuildOpenACCAsteriskSizeExpr(SourceLocation AsteriskLoc); + ExprResult ActOnOpenACCAsteriskSizeExpr(SourceLocation AsteriskLoc); /// Helper type to restore the state of various 'loop' constructs when we run /// into a loop (for, etc) inside the construct. diff --git a/clang/include/clang/Serialization/ASTBitCodes.h b/clang/include/clang/Serialization/ASTBitCodes.h index 1af1f4a10db290..4b79d4b7711905 100644 --- a/clang/include/clang/Serialization/ASTBitCodes.h +++ b/clang/include/clang/Serialization/ASTBitCodes.h @@ -2002,12 +2002,14 @@ enum StmtCode { // SYCLUniqueStableNameExpr EXPR_SYCL_UNIQUE_STABLE_NAME, - // OpenACC Constructs + // OpenACC Constructs/Exprs STMT_OPENACC_COMPUTE_CONSTRUCT, STMT_OPENACC_LOOP_CONSTRUCT, + EXPR_OPENACC_ASTERISK_SIZE, // HLSL Constructs EXPR_HLSL_OUT_ARG, + }; /// The kinds of designators that can occur in a diff --git a/clang/lib/AST/ComputeDependence.cpp b/clang/lib/AST/ComputeDependence.cpp index 6ef49790481aca..8c79df8317a2ec 100644 --- a/clang/lib/AST/ComputeDependence.cpp +++ b/clang/lib/AST/ComputeDependence.cpp @@ -958,3 +958,9 @@ ExprDependence clang::computeDependence(ObjCMessageExpr *E) { D |= A->getDependence(); return D; } + +ExprDependence clang::computeDependence(OpenACCAsteriskSizeExpr *E) { + // This represents a simple asterisk as typed, so cannot be dependent in any + // way. + return ExprDependence::None; +} diff --git a/clang/lib/AST/Expr.cpp b/clang/lib/AST/Expr.cpp index 2e463fc00c6b68..9ecbf121e3fc0d 100644 --- a/clang/lib/AST/Expr.cpp +++ b/clang/lib/AST/Expr.cpp @@ -3640,6 +3640,7 @@ bool Expr::HasSideEffects(const ASTContext &Ctx, case SYCLUniqueStableNameExprClass: case PackIndexingExprClass: case HLSLOutArgExprClass: + case OpenACCAsteriskSizeExprClass: // These never have a side-effect. return false; @@ -5408,3 +5409,13 @@ HLSLOutArgExpr *HLSLOutArgExpr::Create(const ASTContext &C, QualType Ty, HLSLOutArgExpr *HLSLOutArgExpr::CreateEmpty(const ASTContext &C) { return new (C) HLSLOutArgExpr(EmptyShell()); } + +OpenACCAsteriskSizeExpr *OpenACCAsteriskSizeExpr::Create(const ASTContext &C, + SourceLocation Loc) { + return new (C) OpenACCAsteriskSizeExpr(Loc, C.IntTy); +} + +OpenACCAsteriskSizeExpr * +OpenACCAsteriskSizeExpr::CreateEmpty(const ASTContext &C) { + return new (C) OpenACCAsteriskSizeExpr({}, C.IntTy); +} diff --git a/clang/lib/AST/ExprClassification.cpp b/clang/lib/AST/ExprClassification.cpp index 9d97633309ada2..3f37d06cc8f3a0 100644 --- a/clang/lib/AST/ExprClassification.cpp +++ b/clang/lib/AST/ExprClassification.cpp @@ -471,6 +471,7 @@ static Cl::Kinds ClassifyInternal(ASTContext &Ctx, const Expr *E) { case Expr::CoyieldExprClass: return ClassifyInternal(Ctx, cast<CoroutineSuspendExpr>(E)->getResumeExpr()); case Expr::SYCLUniqueStableNameExprClass: + case Expr::OpenACCAsteriskSizeExprClass: return Cl::CL_PRValue; break; diff --git a/clang/lib/AST/ExprConstant.cpp b/clang/lib/AST/ExprConstant.cpp index 3a73cea97fcc32..4d5af96093cfeb 100644 --- a/clang/lib/AST/ExprConstant.cpp +++ b/clang/lib/AST/ExprConstant.cpp @@ -11873,6 +11873,13 @@ class IntExprEvaluator return Success(E->getValue(), E); } + bool VisitOpenACCAsteriskSizeExpr(const OpenACCAsteriskSizeExpr *E) { + // This should not be evaluated during constant expr evaluation, as it + // should always be in an unevaluated context (the args list of a 'gang' or + // 'tile' clause). + return Error(E); + } + bool VisitUnaryReal(const UnaryOperator *E); bool VisitUnaryImag(const UnaryOperator *E); @@ -16908,6 +16915,7 @@ static ICEDiag CheckICE(const Expr* E, const ASTContext &Ctx) { case Expr::GNUNullExprClass: case Expr::SourceLocExprClass: case Expr::EmbedExprClass: + case Expr::OpenACCAsteriskSizeExprClass: return NoDiag(); case Expr::PackIndexingExprClass: diff --git a/clang/lib/AST/ItaniumMangle.cpp b/clang/lib/AST/ItaniumMangle.cpp index 70ed9cb0736f59..769a863c2b6764 100644 --- a/clang/lib/AST/ItaniumMangle.cpp +++ b/clang/lib/AST/ItaniumMangle.cpp @@ -5745,6 +5745,15 @@ void CXXNameMangler::mangleExpression(const Expr *E, unsigned Arity, case Expr::HLSLOutArgExprClass: llvm_unreachable( "cannot mangle hlsl temporary value; mangling wrong thing?"); + case Expr::OpenACCAsteriskSizeExprClass: { + // We shouldn't ever be able to get here, but diagnose anyway. + DiagnosticsEngine &Diags = Context.getDiags(); + unsigned DiagID = Diags.getCustomDiagID( + DiagnosticsEngine::Error, + "cannot yet mangle OpenACC Asterisk Size expression"); + Diags.Report(DiagID); + return; + } } if (AsTemplateArg && !IsPrimaryExpr) diff --git a/clang/lib/AST/JSONNodeDumper.cpp b/clang/lib/AST/JSONNodeDumper.cpp index 565f1e05710c87..ddbe2136a671f3 100644 --- a/clang/lib/AST/JSONNodeDumper.cpp +++ b/clang/lib/AST/JSONNodeDumper.cpp @@ -1354,6 +1354,9 @@ void JSONNodeDumper::VisitSYCLUniqueStableNameExpr( createQualType(E->getTypeSourceInfo()->getType())); } +void JSONNodeDumper::VisitOpenACCAsteriskSizeExpr( + const OpenACCAsteriskSizeExpr *E) {} + void JSONNodeDumper::VisitPredefinedExpr(const PredefinedExpr *PE) { JOS.attribute("name", PredefinedExpr::getIdentKindName(PE->getIdentKind())); } diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp index d864ded33e8d1f..0b34ed6189593e 100644 --- a/clang/lib/AST/OpenACCClause.cpp +++ b/clang/lib/AST/OpenACCClause.cpp @@ -24,6 +24,7 @@ bool OpenACCClauseWithParams::classof(const OpenACCClause *C) { } bool OpenACCClauseWithExprs::classof(const OpenACCClause *C) { return OpenACCWaitClause::classof(C) || OpenACCNumGangsClause::classof(C) || + OpenACCTileClause::classof(C) || OpenACCClauseWithSingleIntExpr::classof(C) || OpenACCClauseWithVarList::classof(C); } @@ -221,6 +222,16 @@ OpenACCNumGangsClause *OpenACCNumGangsClause::Create(const ASTContext &C, return new (Mem) OpenACCNumGangsClause(BeginLoc, LParenLoc, IntExprs, EndLoc); } +OpenACCTileClause *OpenACCTileClause::Create(const ASTContext &C, + SourceLocation BeginLoc, + SourceLocation LParenLoc, + ArrayRef<Expr *> SizeExprs, + SourceLocation EndLoc) { + void *Mem = + C.Allocate(OpenACCTileClause::totalSizeToAlloc<Expr *>(SizeExprs.size())); + return new (Mem) OpenACCTileClause(BeginLoc, LParenLoc, SizeExprs, EndLoc); +} + OpenACCPrivateClause *OpenACCPrivateClause::Create(const ASTContext &C, SourceLocation BeginLoc, SourceLocation LParenLoc, @@ -420,6 +431,13 @@ void OpenACCClausePrinter::VisitNumGangsClause(const OpenACCNumGangsClause &C) { OS << ")"; } +void OpenACCClausePrinter::VisitTileClause(const OpenACCTileClause &C) { + OS << "tile("; + llvm::interleaveComma(C.getSizeExprs(), OS, + [&](const Expr *E) { printExpr(E); }); + OS << ")"; +} + void OpenACCClausePrinter::VisitNumWorkersClause( const OpenACCNumWorkersClause &C) { OS << "num_workers("; diff --git a/clang/lib/AST/StmtPrinter.cpp b/clang/lib/AST/StmtPrinter.cpp index e1b5bec7a50d0a..641f7b52de6dfb 100644 --- a/clang/lib/AST/StmtPrinter.cpp +++ b/clang/lib/AST/StmtPrinter.cpp @@ -1309,6 +1309,10 @@ void StmtPrinter::VisitPredefinedExpr(PredefinedExpr *Node) { OS << PredefinedExpr::getIdentKindName(Node->getIdentKind()); } +void StmtPrinter::VisitOpenACCAsteriskSizeExpr(OpenACCAsteriskSizeExpr *Node) { + OS << '*'; +} + void StmtPrinter::VisitCharacterLiteral(CharacterLiteral *Node) { CharacterLiteral::print(Node->getValue(), Node->getKind(), OS); } diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp index c3812844ab8a31..299ac005c7fdb9 100644 --- a/clang/lib/AST/StmtProfile.cpp +++ b/clang/lib/AST/StmtProfile.cpp @@ -1360,6 +1360,11 @@ void StmtProfiler::VisitPredefinedExpr(const PredefinedExpr *S) { ID.AddInteger(llvm::to_underlying(S->getIdentKind())); } +void StmtProfiler::VisitOpenACCAsteriskSizeExpr( + const OpenACCAsteriskSizeExpr *S) { + VisitExpr(S); +} + void StmtProfiler::VisitIntegerLiteral(const IntegerLiteral *S) { VisitExpr(S); S->getValue().Profile(ID); @@ -2552,6 +2557,11 @@ void OpenACCClauseProfiler::VisitNumGangsClause( Profiler.VisitStmt(E); } +void OpenACCClauseProfiler::VisitTileClause(const OpenACCTileClause &Clause) { + for (auto *E : Clause.getSizeExprs()) + 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 8a74159c7c93e5..15b23d60c3ffab 100644 --- a/clang/lib/AST/TextNodeDumper.cpp +++ b/clang/lib/AST/TextNodeDumper.cpp @@ -381,6 +381,11 @@ void TextNodeDumper::Visit(const OMPClause *C) { OS << " <implicit>"; } +void TextNodeDumper::VisitOpenACCAsteriskSizeExpr( + const OpenACCAsteriskSizeExpr *E) { + // Nothing to do here, only location exists, and that is printed elsewhere. +} + void TextNodeDumper::Visit(const OpenACCClause *C) { if (!C) { ColorScope Color(OS, ShowColors, NullColor); @@ -414,6 +419,7 @@ void TextNodeDumper::Visit(const OpenACCClause *C) { case OpenACCClauseKind::Private: case OpenACCClauseKind::Self: case OpenACCClauseKind::Seq: + case OpenACCClauseKind::Tile: case OpenACCClauseKind::VectorLength: // The condition expression will be printed as a part of the 'children', // but print 'clause' here so it is clear what is happening from the dump. diff --git a/clang/lib/CodeGen/CGExprScalar.cpp b/clang/lib/CodeGen/CGExprScalar.cpp index b7f5b932c56b6f..7529d4465f2c34 100644 --- a/clang/lib/CodeGen/CGExprScalar.cpp +++ b/clang/lib/CodeGen/CGExprScalar.cpp @@ -531,6 +531,10 @@ class ScalarExprEmitter return CGF.getOrCreateOpaqueRValueMapping(E).getScalarVal(); } + Value *VisitOpenACCAsteriskSizeExpr(OpenACCAsteriskSizeExpr *E) { + llvm_unreachable("Codegen for this isn't defined/implemented"); + } + // l-values. Value *VisitDeclRefExpr(DeclRefExpr *E) { if (CodeGenFunction::ConstantEmission Constant = CGF.tryEmitAsConstant(E)) diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp index 87673beb34300e..b27e50b147f4a8 100644 --- a/clang/lib/Parse/ParseOpenACC.cpp +++ b/clang/lib/Parse/ParseOpenACC.cpp @@ -743,38 +743,50 @@ bool Parser::ParseOpenACCDeviceTypeList( // int-expr // Note that this is specified under 'gang-arg-list', but also applies to 'tile' // via reference. -bool Parser::ParseOpenACCSizeExpr() { - // FIXME: Ensure these are constant expressions. - +ExprResult Parser::ParseOpenACCSizeExpr(OpenACCClauseKind CK) { // The size-expr ends up being ambiguous when only looking at the current // token, as it could be a deref of a variable/expression. if (getCurToken().is(tok::star) && NextToken().isOneOf(tok::comma, tok::r_paren, tok::annot_pragma_openacc_end)) { - ConsumeToken(); - return false; + SourceLocation AsteriskLoc = ConsumeToken(); + return getActions().OpenACC().ActOnOpenACCAsteriskSizeExpr(AsteriskLoc); } - return getActions() - .CorrectDelayedTyposInExpr(ParseAssignmentExpression()) - .isInvalid(); + ExprResult SizeExpr = + getActions().CorrectDelayedTyposInExpr(ParseConstantExpression()); + + if (!SizeExpr.isUsable()) + return SizeExpr; + + SizeExpr = getActions().OpenACC().ActOnIntExpr( + OpenACCDirectiveKind::Invalid, CK, SizeExpr.get()->getBeginLoc(), + SizeExpr.get()); + + return SizeExpr; } -bool Parser::ParseOpenACCSizeExprList() { - if (ParseOpenACCSizeExpr()) { +bool Parser::ParseOpenACCSizeExprList( + OpenACCClauseKind CK, llvm::SmallVectorImpl<Expr *> &SizeExprs) { + ExprResult SizeExpr = ParseOpenACCSizeExpr(CK); + if (!SizeExpr.isUsable()) { SkipUntil(tok::r_paren, tok::annot_pragma_openacc_end, Parser::StopBeforeMatch); - return false; + return true; } + SizeExprs.push_back(SizeExpr.get()); + while (!getCurToken().isOneOf(tok::r_paren, tok::annot_pragma_openacc_end)) { ExpectAndConsume(tok::comma); - if (ParseOpenACCSizeExpr()) { + SizeExpr = ParseOpenACCSizeExpr(CK); + if (!SizeExpr.isUsable()) { SkipUntil(tok::r_paren, tok::annot_pragma_openacc_end, Parser::StopBeforeMatch); - return false; + return true; } + SizeExprs.push_back(SizeExpr.get()); } return false; } @@ -792,7 +804,7 @@ bool Parser::ParseOpenACCGangArg(SourceLocation GangLoc) { // 'static' just takes a size-expr, which is an int-expr or an asterisk. ConsumeToken(); ConsumeToken(); - return ParseOpenACCSizeExpr(); + return ParseOpenACCSizeExpr(OpenACCClauseKind::Gang).isInvalid(); } if (isOpenACCSpecialToken(OpenACCSpecialTokenKind::Dim, getCurToken()) && @@ -1052,12 +1064,16 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams( } break; } - case OpenACCClauseKind::Tile: - if (ParseOpenACCSizeExprList()) { + case OpenACCClauseKind::Tile: { + llvm::SmallVector<Expr *> SizeExprs; + if (ParseOpenACCSizeExprList(OpenACCClauseKind::Tile, SizeExprs)) { Parens.skipToEnd(); return OpenACCCanContinue(); } + + ParsedClause.setIntExprDetails(std::move(SizeExprs)); break; + } default: llvm_unreachable("Not a required parens type?"); } diff --git a/clang/lib/Sema/SemaExceptionSpec.cpp b/clang/lib/Sema/SemaExceptionSpec.cpp index 8aedbfcf878a11..dbddd6c370aa07 100644 --- a/clang/lib/Sema/SemaExceptionSpec.cpp +++ b/clang/lib/Sema/SemaExceptionSpec.cpp @@ -1593,6 +1593,8 @@ CanThrowResult Sema::canThrow(const Stmt *S) { case Stmt::SYCLUniqueStableNameExprClass: return CT_Cannot; + case Stmt::OpenACCAsteriskSizeExprClass: + return CT_Cannot; case Stmt::NoStmtClass: llvm_unreachable("Invalid class for statement"); } diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp index dfaf726891cfc8..f3840c6ffba82b 100644 --- a/clang/lib/Sema/SemaOpenACC.cpp +++ b/clang/lib/Sema/SemaOpenACC.cpp @@ -354,6 +354,17 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind, return false; } } + case OpenACCClauseKind::Tile: { + switch (DirectiveKind) { + case OpenACCDirectiveKind::Loop: + case OpenACCDirectiveKind::ParallelLoop: + case OpenACCDirectiveKind::SerialLoop: + case OpenACCDirectiveKind::KernelsLoop: + return true; + default: + return false; + } + } default: // Do nothing so we can go to the 'unimplemented' diagnostic instead. @@ -526,6 +537,36 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitDefaultClause( Clause.getLParenLoc(), Clause.getEndLoc()); } +OpenACCClause *SemaOpenACCClauseVisitor::VisitTileClause( + SemaOpenACC::OpenACCParsedClause &Clause) { + if (Clause.getDirectiveKind() != OpenACCDirectiveKind::Loop) + return isNotImplemented(); + + // Duplicates here are not really sensible. We could possible permit + // multiples if they all had the same value, but there isn't really a good + // reason to do so. Also, this simplifies the suppression of duplicates, in + // that we know if we 'find' one after instantiation, that it is the same + // clause, which simplifies instantiation/checking/etc. + if (checkAlreadyHasClauseOfKind(SemaRef, ExistingClauses, Clause)) + return nullptr; + + llvm::SmallVector<Expr *> NewSizeExprs; + + // Make sure these are all positive constant expressions or *. + for (Expr *E : Clause.getIntExprs()) { + ExprResult Res = SemaRef.CheckTileSizeExpr(E); + + if (!Res.isUsable()) + return nullptr; + + NewSizeExprs.push_back(Res.get()); + } + + return OpenACCTileClause::Create(Ctx, Clause.getBeginLoc(), + Clause.getLParenLoc(), NewSizeExprs, + Clause.getEndLoc()); +} + OpenACCClause *SemaOpenACCClauseVisitor::VisitIfClause( SemaOpenACC::OpenACCParsedClause &Clause) { // Restrictions only properly implemented on 'compute' constructs, and @@ -1698,6 +1739,34 @@ ExprResult SemaOpenACC::CheckCollapseLoopCount(Expr *LoopCount) { ConstantExpr::Create(getASTContext(), LoopCount, APValue{*ICE})}; } +ExprResult SemaOpenACC::CheckTileSizeExpr(Expr *SizeExpr) { + if (!SizeExpr) + return ExprError(); + + assert((SizeExpr->isInstantiationDependent() || + SizeExpr->getType()->isIntegerType()) && + "size argument non integer?"); + + // If dependent, or an asterisk, the expression is fine. + if (SizeExpr->isInstantiationDependent() || + isa<OpenACCAsteriskSizeExpr>(SizeExpr)) + return ExprResult{SizeExpr}; + + std::optional<llvm::APSInt> ICE = + SizeExpr->getIntegerConstantExpr(getASTContext()); + + // OpenACC 3.3 2.9.8 + // where each tile size is a constant positive integer expression or asterisk. + if (!ICE || *ICE <= 0) { + Diag(SizeExpr->getBeginLoc(), diag::err_acc_size_expr_value) + << ICE.has_value() << ICE.value_or(llvm::APSInt{}).getExtValue(); + return ExprError(); + } + + return ExprResult{ + ConstantExpr::Create(getASTContext(), SizeExpr, APValue{*ICE})}; +} + void SemaOpenACC::ActOnWhileStmt(SourceLocation WhileLoc) { if (!getLangOpts().OpenACC) return; @@ -1944,3 +2013,13 @@ bool SemaOpenACC::ActOnStartDeclDirective(OpenACCDirectiveKind K, } DeclGroupRef SemaOpenACC::ActOnEndDeclDirective() { return DeclGroupRef{}; } + +ExprResult +SemaOpenACC::BuildOpenACCAsteriskSizeExpr(SourceLocation AsteriskLoc) { + return OpenACCAsteriskSizeExpr::Create(getASTContext(), AsteriskLoc); +} + +ExprResult +SemaOpenACC::ActOnOpenACCAsteriskSizeExpr(SourceLocation AsteriskLoc) { + return BuildOpenACCAsteriskSizeExpr(AsteriskLoc); +} diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h index 04308b2b57d8a0..76efc9fe831854 100644 --- a/clang/lib/Sema/TreeTransform.h +++ b/clang/lib/Sema/TreeTransform.h @@ -4090,6 +4090,10 @@ class TreeTransform { OpenACCDirectiveKind::Loop, BeginLoc, DirLoc, EndLoc, Clauses, Loop); } + ExprResult RebuildOpenACCAsteriskSizeExpr(SourceLocation AsteriskLoc) { + return getSema().OpenACC().ActOnOpenACCAsteriskSizeExpr(AsteriskLoc); + } + private: TypeLoc TransformTypeInObjectScope(TypeLoc TL, QualType ObjectType, @@ -11871,6 +11875,36 @@ void OpenACCClauseTransform<Derived>::VisitCollapseClause( ParsedClause.getLParenLoc(), ParsedClause.isForce(), ParsedClause.getLoopCount(), ParsedClause.getEndLoc()); } + +template <typename Derived> +void OpenACCClauseTransform<Derived>::VisitTileClause( + const OpenACCTileClause &C) { + + llvm::SmallVector<Expr *> TransformedExprs; + + for (Expr *E : C.getSizeExprs()) { + ExprResult NewSizeExpr = Self.TransformExpr(E); + + if (!NewSizeExpr.isUsable()) + return; + + NewSizeExpr = Self.getSema().OpenACC().ActOnIntExpr( + OpenACCDirectiveKind::Invalid, ParsedClause.getClauseKind(), + NewSizeExpr.get()->getBeginLoc(), NewSizeExpr.get()); + + NewSizeExpr = Self.getSema().OpenACC().CheckTileSizeExpr(NewSizeExpr.get()); + + if (!NewSizeExpr.isUsable()) + return; + TransformedExprs.push_back(NewSizeExpr.get()); + } + + ParsedClause.setIntExprDetails(TransformedExprs); + NewClause = OpenACCTileClause::Create( + Self.getSema().getASTContext(), ParsedClause.getBeginLoc(), + ParsedClause.getLParenLoc(), ParsedClause.getIntExprs(), + ParsedClause.getEndLoc()); +} } // namespace template <typename Derived> OpenACCClause *TreeTransform<Derived>::TransformOpenACCClause( @@ -11957,6 +11991,15 @@ TreeTransform<Derived>::TransformOpenACCLoopConstruct(OpenACCLoopConstruct *C) { TransformedClauses, Loop); } +template <typename Derived> +ExprResult TreeTransform<Derived>::TransformOpenACCAsteriskSizeExpr( + OpenACCAsteriskSizeExpr *E) { + if (getDerived().AlwaysRebuild()) + return getDerived().RebuildOpenACCAsteriskSizeExpr(E->getLocation()); + // Nothing can ever change, so there is never anything to transform. + return E; +} + //===----------------------------------------------------------------------===// // Expression transformation //===----------------------------------------------------------------------===// diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp index 4b599b7dbf3ce8..97b79bd1381c02 100644 --- a/clang/lib/Serialization/ASTReader.cpp +++ b/clang/lib/Serialization/ASTReader.cpp @@ -12306,6 +12306,15 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() { return OpenACCCollapseClause::Create(getContext(), BeginLoc, LParenLoc, HasForce, LoopCount, EndLoc); } + case OpenACCClauseKind::Tile: { + SourceLocation LParenLoc = readSourceLocation(); + unsigned NumClauses = readInt(); + llvm::SmallVector<Expr *> SizeExprs; + for (unsigned I = 0; I < NumClauses; ++I) + SizeExprs.push_back(readSubExpr()); + return OpenACCTileClause::Create(getContext(), BeginLoc, LParenLoc, + SizeExprs, EndLoc); + } case OpenACCClauseKind::Finalize: case OpenACCClauseKind::IfPresent: @@ -12322,7 +12331,6 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() { case OpenACCClauseKind::Bind: case OpenACCClauseKind::DeviceNum: case OpenACCClauseKind::DefaultAsync: - case OpenACCClauseKind::Tile: case OpenACCClauseKind::Gang: case OpenACCClauseKind::Invalid: llvm_unreachable("Clause serialization not yet implemented"); diff --git a/clang/lib/Serialization/ASTReaderStmt.cpp b/clang/lib/Serialization/ASTReaderStmt.cpp index 84743a52d4c8b8..2038fe7829af9b 100644 --- a/clang/lib/Serialization/ASTReaderStmt.cpp +++ b/clang/lib/Serialization/ASTReaderStmt.cpp @@ -575,6 +575,11 @@ void ASTStmtReader::VisitConstantExpr(ConstantExpr *E) { E->setSubExpr(Record.readSubExpr()); } +void ASTStmtReader::VisitOpenACCAsteriskSizeExpr(OpenACCAsteriskSizeExpr *E) { + VisitExpr(E); + E->setAsteriskLocation(readSourceLocation()); +} + void ASTStmtReader::VisitSYCLUniqueStableNameExpr(SYCLUniqueStableNameExpr *E) { VisitExpr(E); @@ -3054,6 +3059,10 @@ Stmt *ASTReader::ReadStmtFromStream(ModuleFile &F) { S = SYCLUniqueStableNameExpr::CreateEmpty(Context); break; + case EXPR_OPENACC_ASTERISK_SIZE: + S = OpenACCAsteriskSizeExpr::CreateEmpty(Context); + break; + case EXPR_PREDEFINED: S = PredefinedExpr::CreateEmpty( Context, diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp index 375ddc90482b27..836532ca402ffc 100644 --- a/clang/lib/Serialization/ASTWriter.cpp +++ b/clang/lib/Serialization/ASTWriter.cpp @@ -8167,6 +8167,14 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) { AddStmt(const_cast<Expr *>(CC->getLoopCount())); return; } + case OpenACCClauseKind::Tile: { + const auto *TC = cast<OpenACCTileClause>(C); + writeSourceLocation(TC->getLParenLoc()); + writeUInt32(TC->getSizeExprs().size()); + for (Expr *E : TC->getSizeExprs()) + AddStmt(E); + return; + } case OpenACCClauseKind::Finalize: case OpenACCClauseKind::IfPresent: @@ -8183,7 +8191,6 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) { case OpenACCClauseKind::Bind: case OpenACCClauseKind::DeviceNum: case OpenACCClauseKind::DefaultAsync: - case OpenACCClauseKind::Tile: case OpenACCClauseKind::Gang: case OpenACCClauseKind::Invalid: llvm_unreachable("Clause serialization not yet implemented"); diff --git a/clang/lib/Serialization/ASTWriterStmt.cpp b/clang/lib/Serialization/ASTWriterStmt.cpp index 837136600181c1..64e4894dc29fb7 100644 --- a/clang/lib/Serialization/ASTWriterStmt.cpp +++ b/clang/lib/Serialization/ASTWriterStmt.cpp @@ -643,6 +643,12 @@ void ASTStmtWriter::VisitConstantExpr(ConstantExpr *E) { Code = serialization::EXPR_CONSTANT; } +void ASTStmtWriter::VisitOpenACCAsteriskSizeExpr(OpenACCAsteriskSizeExpr *E) { + VisitExpr(E); + Record.AddSourceLocation(E->getLocation()); + Code = serialization::EXPR_OPENACC_ASTERISK_SIZE; +} + void ASTStmtWriter::VisitSYCLUniqueStableNameExpr(SYCLUniqueStableNameExpr *E) { VisitExpr(E); diff --git a/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp b/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp index fdabba46992b08..b1919d7027cf4d 100644 --- a/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp +++ b/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp @@ -1966,6 +1966,7 @@ void ExprEngine::Visit(const Stmt *S, ExplodedNode *Pred, case Stmt::OMPArrayShapingExprClass: case Stmt::OMPIteratorExprClass: case Stmt::SYCLUniqueStableNameExprClass: + case Stmt::OpenACCAsteriskSizeExprClass: case Stmt::TypeTraitExprClass: { Bldr.takeNodes(Pred); ExplodedNodeSet preVisit; diff --git a/clang/test/AST/ast-print-openacc-loop-construct.cpp b/clang/test/AST/ast-print-openacc-loop-construct.cpp index ae1f7964f019eb..aee4591cab428f 100644 --- a/clang/test/AST/ast-print-openacc-loop-construct.cpp +++ b/clang/test/AST/ast-print-openacc-loop-construct.cpp @@ -2,6 +2,7 @@ struct SomeStruct{}; +constexpr int get_value() { return 1; } void foo() { // CHECK: #pragma acc loop // CHECK-NEXT: for (;;) @@ -82,4 +83,16 @@ void foo() { #pragma acc loop collapse(force:2) for(;;) for(;;); + +// CHECK: #pragma acc loop tile(1, 3, *, get_value()) +// CHECK-NEXT: for (;;) +// CHECK-NEXT: for (;;) +// CHECK-NEXT: for (;;) +// CHECK-NEXT: for (;;) +// CHECK-NEXT: ; +#pragma acc loop tile(1, 3, *, get_value()) + for(;;) + for(;;) + for(;;) + for(;;); } diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c index 6c9ce4ad5e1969..a752b2ebd18b6a 100644 --- a/clang/test/ParserOpenACC/parse-clauses.c +++ b/clang/test/ParserOpenACC/parse-clauses.c @@ -1153,8 +1153,7 @@ void Tile() { // expected-note@+1{{to match this '('}} #pragma acc loop tile( for(;;){} - // expected-error@+2{{expected expression}} - // expected-warning@+1{{OpenACC clause 'tile' not yet implemented, clause ignored}} + // expected-error@+1{{expected expression}} #pragma acc loop tile() for(;;){} // expected-error@+3{{expected expression}} @@ -1162,41 +1161,32 @@ void Tile() { // expected-note@+1{{to match this '('}} #pragma acc loop tile(, for(;;){} - // expected-error@+2{{expected expression}} - // expected-warning@+1{{OpenACC clause 'tile' not yet implemented, clause ignored}} + // expected-error@+1{{expected expression}} #pragma acc loop tile(,) for(;;){} - // expected-error@+2{{use of undeclared identifier 'invalid'}} - // expected-warning@+1{{OpenACC clause 'tile' not yet implemented, clause ignored}} + // expected-error@+1{{use of undeclared identifier 'invalid'}} #pragma acc loop tile(returns_int(), *, invalid, *) for(;;){} - // expected-error@+2{{expected expression}} - // expected-warning@+1{{OpenACC clause 'tile' not yet implemented, clause ignored}} + // expected-error@+1{{expected expression}} #pragma acc loop tile(returns_int() *, Foo, *) for(;;){} - // expected-error@+2{{indirection requires pointer operand ('int' invalid)}} - // expected-warning@+1{{OpenACC clause 'tile' not yet implemented, clause ignored}} + // expected-error@+1{{indirection requires pointer operand ('int' invalid)}} #pragma acc loop tile(* returns_int() , *) for(;;){} - // expected-warning@+1{{OpenACC clause 'tile' not yet implemented, clause ignored}} #pragma acc loop tile(*) for(;;){} - // expected-warning@+1{{OpenACC clause 'tile' not yet implemented, clause ignored}} + // expected-error@+1{{OpenACC 'tile' clause size expression must be an asterisk or a constant expression}} #pragma acc loop tile(*Foo, *Foo) for(;;){} - // expected-warning@+1{{OpenACC clause 'tile' not yet implemented, clause ignored}} #pragma acc loop tile(5) for(;;){} - // expected-warning@+1{{OpenACC clause 'tile' not yet implemented, clause ignored}} #pragma acc loop tile(*, 5) for(;;){} - // expected-warning@+1{{OpenACC clause 'tile' not yet implemented, clause ignored}} #pragma acc loop tile(5, *) for(;;){} - // expected-warning@+1{{OpenACC clause 'tile' not yet implemented, clause ignored}} #pragma acc loop tile(5, *, 3, *) for(;;){} } diff --git a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c index 26f0315fb86f1a..d08497a7782edb 100644 --- a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c +++ b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c @@ -211,8 +211,7 @@ void uses() { while(1); #pragma acc kernels device_type(*) async while(1); - // expected-error@+2{{OpenACC clause 'tile' may not follow a 'device_type' clause in a compute construct}} - // expected-note@+1{{previous clause is here}} + // 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}} 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 3212c19d089fc9..f66ce6991acb1a 100644 --- a/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c +++ b/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c @@ -165,8 +165,7 @@ void uses() { // expected-error@+1{{OpenACC 'async' clause is not valid on 'loop' directive}} #pragma acc loop auto async for(;;); - // expected-warning@+1{{OpenACC clause 'tile' not yet implemented}} -#pragma acc loop auto tile(Var, 1) +#pragma acc loop auto tile(1+2, 1) for(;;); // expected-warning@+1{{OpenACC clause 'gang' not yet implemented}} #pragma acc loop auto gang @@ -303,8 +302,7 @@ void uses() { // expected-error@+1{{OpenACC 'async' clause is not valid on 'loop' directive}} #pragma acc loop async auto for(;;); - // expected-warning@+1{{OpenACC clause 'tile' not yet implemented}} -#pragma acc loop tile(Var, 1) auto +#pragma acc loop tile(1+2, 1) auto for(;;); // expected-warning@+1{{OpenACC clause 'gang' not yet implemented}} #pragma acc loop gang auto @@ -442,8 +440,7 @@ void uses() { // expected-error@+1{{OpenACC 'async' clause is not valid on 'loop' directive}} #pragma acc loop independent async for(;;); - // expected-warning@+1{{OpenACC clause 'tile' not yet implemented}} -#pragma acc loop independent tile(Var, 1) +#pragma acc loop independent tile(1+2, 1) for(;;); // expected-warning@+1{{OpenACC clause 'gang' not yet implemented}} #pragma acc loop independent gang @@ -580,8 +577,7 @@ void uses() { // expected-error@+1{{OpenACC 'async' clause is not valid on 'loop' directive}} #pragma acc loop async independent for(;;); - // expected-warning@+1{{OpenACC clause 'tile' not yet implemented}} -#pragma acc loop tile(Var, 1) independent +#pragma acc loop tile(1+2, 1) independent for(;;); // expected-warning@+1{{OpenACC clause 'gang' not yet implemented}} #pragma acc loop gang independent @@ -728,8 +724,7 @@ void uses() { // expected-error@+1{{OpenACC 'async' clause is not valid on 'loop' directive}} #pragma acc loop seq async for(;;); - // expected-warning@+1{{OpenACC clause 'tile' not yet implemented}} -#pragma acc loop seq tile(Var, 1) +#pragma acc loop seq tile(1+2, 1) for(;;); // expected-error@+1{{OpenACC 'wait' clause is not valid on 'loop' directive}} #pragma acc loop seq wait @@ -875,8 +870,7 @@ void uses() { // expected-error@+1{{OpenACC 'async' clause is not valid on 'loop' directive}} #pragma acc loop async seq for(;;); - // expected-warning@+1{{OpenACC clause 'tile' not yet implemented}} -#pragma acc loop tile(Var, 1) seq +#pragma acc loop tile(1+2, 1) seq for(;;); // expected-error@+1{{OpenACC 'wait' clause is not valid on 'loop' directive}} #pragma acc loop wait seq diff --git a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c index 47c9239f4f0e96..aa7c3201c8319c 100644 --- a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c +++ b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c @@ -188,8 +188,8 @@ void uses() { // expected-error@+1{{OpenACC 'async' clause is not valid on 'loop' directive}} #pragma acc loop device_type(*) async for(;;); - // expected-warning@+1{{OpenACC clause 'tile' not yet implemented, clause ignored}} -#pragma acc loop device_type(*) tile(Var, 1) + +#pragma acc loop device_type(*) tile(*, 1) for(;;); // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}} #pragma acc loop dtype(*) gang diff --git a/clang/test/SemaOpenACC/loop-construct-tile-ast.cpp b/clang/test/SemaOpenACC/loop-construct-tile-ast.cpp new file mode 100644 index 00000000000000..977c7e3b6d4cf5 --- /dev/null +++ b/clang/test/SemaOpenACC/loop-construct-tile-ast.cpp @@ -0,0 +1,231 @@ +// 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 + +struct S { + constexpr S(){}; + constexpr operator auto() {return 1;} +}; + +void NormalUses() { + // CHECK: FunctionDecl{{.*}}NormalUses + // CHECK-NEXT: CompoundStmt + +#pragma acc loop tile(S{}, 1, 2, *) + for(;;) + for(;;) + for(;;) + for(;;); + // CHECK-NEXT: OpenACCLoopConstruct + // CHECK-NEXT: tile clause + // CHECK-NEXT: ConstantExpr{{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <UserDefinedConversion> + // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int' + // CHECK-NEXT: MemberExpr{{.*}} .operator auto + // CHECK-NEXT: MaterializeTemporaryExpr{{.*}} 'S' lvalue + // CHECK-NEXT: CXXTemporaryObjectExpr{{.*}} 'S' 'void ()' list + // CHECK-NEXT: ConstantExpr{{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1 + // CHECK-NEXT: ConstantExpr{{.*}} 'int' + // CHECK-NEXT: value: Int 2 + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 2 + // CHECK-NEXT: OpenACCAsteriskSizeExpr{{.*}}'int' + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: NullStmt +} + +template<typename T, unsigned Value> +void TemplUses() { + // CHECK: FunctionTemplateDecl{{.*}}TemplUses + // CHECK-NEXT: TemplateTypeParmDecl {{.*}} referenced typename depth 0 index 0 T + // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} referenced 'unsigned int' depth 0 index 1 Value + // CHECK-NEXT: FunctionDecl{{.*}} TemplUses 'void ()' + // CHECK-NEXT: CompoundStmt + +#pragma acc loop tile(S{}, T{}, *, T{} + S{}, Value, Value + 3) + for(;;) + for(;;) + for(;;) + for(;;) + for(;;) + for(;;); + // CHECK-NEXT: OpenACCLoopConstruct + // CHECK-NEXT: tile clause + // + // CHECK-NEXT: ConstantExpr{{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <UserDefinedConversion> + // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int' + // CHECK-NEXT: MemberExpr{{.*}} .operator auto + // CHECK-NEXT: MaterializeTemporaryExpr{{.*}} 'S' lvalue + // CHECK-NEXT: CXXTemporaryObjectExpr{{.*}} 'S' 'void ()' list + // + // CHECK-NEXT: CXXUnresolvedConstructExpr{{.*}} 'T' 'T' list + // CHECK-NEXT: InitListExpr{{.*}}'void' + // + // CHECK-NEXT: OpenACCAsteriskSizeExpr{{.*}}'int' + // + // CHECK-NEXT: BinaryOperator{{.*}}'<dependent type>' '+' + // CHECK-NEXT: CXXUnresolvedConstructExpr{{.*}} 'T' 'T' list + // CHECK-NEXT: InitListExpr{{.*}}'void' + // CHECK-NEXT: CXXTemporaryObjectExpr{{.*}} 'S' 'void ()' list + // + // CHECK-NEXT: DeclRefExpr{{.*}}'unsigned int' NonTypeTemplateParm{{.*}}'Value' 'unsigned int' + // + // CHECK-NEXT: BinaryOperator{{.*}}'unsigned int' '+' + // CHECK-NEXT: DeclRefExpr{{.*}}'unsigned int' NonTypeTemplateParm{{.*}}'Value' 'unsigned int' + // CHECK-NEXT: ImplicitCastExpr{{.*}} 'unsigned int' <IntegralCast> + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 3 + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: NullStmt + + // Instantiation: + // CHECK-NEXT: FunctionDecl{{.*}} used TemplUses 'void ()' implicit_instantiation + // CHECK-NEXT: TemplateArgument type 'S' + // CHECK-NEXT: RecordType{{.*}} 'S' + // CHECK-NEXT: CXXRecord{{.*}} 'S' + // CHECK-NEXT: TemplateArgument integral '2U' + // CHECK-NEXT: CompoundStmt + + // CHECK-NEXT: OpenACCLoopConstruct + // CHECK-NEXT: tile clause + // + // CHECK-NEXT: ConstantExpr{{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <UserDefinedConversion> + // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int' + // CHECK-NEXT: MemberExpr{{.*}} .operator auto + // CHECK-NEXT: MaterializeTemporaryExpr{{.*}} 'S' lvalue + // CHECK-NEXT: CXXTemporaryObjectExpr{{.*}} 'S' 'void ()' list + // + // CHECK-NEXT: ConstantExpr{{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <UserDefinedConversion> + // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int' + // CHECK-NEXT: MemberExpr{{.*}} .operator auto + // CHECK-NEXT: MaterializeTemporaryExpr{{.*}} 'S' lvalue + // CHECK-NEXT: CXXTemporaryObjectExpr{{.*}} 'S' 'void ()' list + // + // CHECK-NEXT: OpenACCAsteriskSizeExpr{{.*}}'int' + // + // CHECK-NEXT: ConstantExpr{{.*}} 'int' + // CHECK-NEXT: value: Int 2 + // CHECK-NEXT: BinaryOperator{{.*}}'int' '+' + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <UserDefinedConversion> + // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int' + // CHECK-NEXT: MemberExpr{{.*}} .operator auto + // CHECK-NEXT: MaterializeTemporaryExpr{{.*}} 'S' lvalue + // CHECK-NEXT: CXXTemporaryObjectExpr{{.*}} 'S' 'void ()' list + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <UserDefinedConversion> + // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int' + // CHECK-NEXT: MemberExpr{{.*}} .operator auto + // CHECK-NEXT: MaterializeTemporaryExpr{{.*}} 'S' lvalue + // CHECK-NEXT: CXXTemporaryObjectExpr{{.*}} 'S' 'void ()' list + // + // CHECK-NEXT: ConstantExpr{{.*}} 'unsigned int' + // CHECK-NEXT: value: Int 2 + // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}} 'unsigned int' + // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'unsigned int' depth 0 index 1 Value + // CHECK-NEXT: IntegerLiteral{{.*}} 'unsigned int' 2 + // + // CHECK-NEXT: ConstantExpr{{.*}} 'unsigned int' + // CHECK-NEXT: value: Int 5 + // CHECK-NEXT: BinaryOperator{{.*}}'unsigned int' '+' + // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}} 'unsigned int' + // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'unsigned int' depth 0 index 1 Value + // CHECK-NEXT: IntegerLiteral{{.*}} 'unsigned int' 2 + // CHECK-NEXT: ImplicitCastExpr{{.*}} 'unsigned int' <IntegralCast> + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 3 + + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: NullStmt + +} + +void Inst() { + TemplUses<S, 2>(); +} +#endif // PCH_HELPER diff --git a/clang/test/SemaOpenACC/loop-construct-tile-clause.cpp b/clang/test/SemaOpenACC/loop-construct-tile-clause.cpp new file mode 100644 index 00000000000000..28eadde7416c35 --- /dev/null +++ b/clang/test/SemaOpenACC/loop-construct-tile-clause.cpp @@ -0,0 +1,95 @@ +// RUN: %clang_cc1 %s -fopenacc -verify + +constexpr int three() { return 3; } +constexpr int one() { return 1; } +constexpr int neg() { return -1; } +constexpr int zero() { return 0; } + +struct NotConstexpr { + constexpr NotConstexpr(){}; + + operator int(){ return 1; } +}; +struct ConvertsNegative { + constexpr ConvertsNegative(){}; + + constexpr operator int(){ return -1; } +}; +struct ConvertsOne{ + constexpr ConvertsOne(){}; + + constexpr operator int(){ return 1; } +}; + +struct ConvertsThree{ + constexpr ConvertsThree(){}; + + constexpr operator int(){ return 3; } +}; + +template<typename T, int Val> +void negative_zero_constexpr_templ() { + // expected-error@+1 2{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to 0}} +#pragma acc loop tile(*, T{}) + for(;;) + for(;;); + + // expected-error@+1{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to -1}} +#pragma acc loop tile(Val, *) + for(;;) + for(;;); + + // expected-error@+1{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to 0}} +#pragma acc loop tile(zero(), *) + for(;;) + for(;;); +} + +void negative_zero_constexpr() { + negative_zero_constexpr_templ<int, 1>(); // expected-note{{in instantiation of function template specialization}} + negative_zero_constexpr_templ<int, -1>(); // expected-note{{in instantiation of function template specialization}} + + // expected-error@+1{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to 0}} +#pragma acc loop tile(0, *) + for(;;) + for(;;); + + // expected-error@+1{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to 0}} +#pragma acc loop tile(1, 0) + for(;;) + for(;;); + + // expected-error@+1{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to -1}} +#pragma acc loop tile(1, -1) + for(;;) + for(;;); + + // expected-error@+1{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to -1}} +#pragma acc loop tile(-1, 0) + for(;;) + for(;;); + + // expected-error@+1{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to 0}} +#pragma acc loop tile(zero(), 0) + for(;;) + for(;;); + + // expected-error@+1{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to -1}} +#pragma acc loop tile(1, neg()) + for(;;) + for(;;); + + // expected-error@+1{{OpenACC 'tile' clause size expression must be an asterisk or a constant expression}} +#pragma acc loop tile(NotConstexpr{}) + for(;;); + + // expected-error@+1{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to -1}} +#pragma acc loop tile(1, ConvertsNegative{}) + for(;;) + for(;;); + +#pragma acc loop tile(*, ConvertsOne{}) + for(;;) + for(;;); + +} diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp index d188f794bad20e..e6c323775c999e 100644 --- a/clang/tools/libclang/CIndex.cpp +++ b/clang/tools/libclang/CIndex.cpp @@ -2839,6 +2839,11 @@ void OpenACCClauseEnqueue::VisitNumGangsClause(const OpenACCNumGangsClause &C) { Visitor.AddStmt(IE); } +void OpenACCClauseEnqueue::VisitTileClause(const OpenACCTileClause &C) { + for (Expr *IE : C.getSizeExprs()) + Visitor.AddStmt(IE); +} + void OpenACCClauseEnqueue::VisitPrivateClause(const OpenACCPrivateClause &C) { VisitVarList(C); } diff --git a/clang/tools/libclang/CXCursor.cpp b/clang/tools/libclang/CXCursor.cpp index 4e068f272a153f..562228cc334f3a 100644 --- a/clang/tools/libclang/CXCursor.cpp +++ b/clang/tools/libclang/CXCursor.cpp @@ -337,6 +337,7 @@ CXCursor cxcursor::MakeCXCursor(const Stmt *S, const Decl *Parent, case Stmt::SYCLUniqueStableNameExprClass: case Stmt::EmbedExprClass: case Stmt::HLSLOutArgExprClass: + case Stmt::OpenACCAsteriskSizeExprClass: K = CXCursor_UnexposedExpr; break; _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
