Author: erichkeane Date: 2025-03-03T13:57:23-08:00 New Revision: d5cec386c14ac46ee252da29f5bd766db0adb6d0
URL: https://github.com/llvm/llvm-project/commit/d5cec386c14ac46ee252da29f5bd766db0adb6d0 DIFF: https://github.com/llvm/llvm-project/commit/d5cec386c14ac46ee252da29f5bd766db0adb6d0.diff LOG: [OpenACC] Implement 'cache' construct AST/Sema This statement level construct takes no clauses and has no associated statement, and simply labels a number of array elements as valid for caching. The implementation here is pretty simple, but it is a touch of a special case for parsing, so the parsing code reflects that. Added: clang/test/AST/ast-print-openacc-cache-construct.cpp clang/test/SemaOpenACC/cache-construct-ast.cpp clang/test/SemaOpenACC/cache-construct.cpp Modified: clang/include/clang-c/Index.h clang/include/clang/AST/RecursiveASTVisitor.h clang/include/clang/AST/StmtOpenACC.h clang/include/clang/AST/TextNodeDumper.h clang/include/clang/Basic/DiagnosticSemaKinds.td 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/StmtOpenACC.cpp clang/lib/AST/StmtPrinter.cpp clang/lib/AST/StmtProfile.cpp clang/lib/AST/TextNodeDumper.cpp clang/lib/CodeGen/CGStmt.cpp clang/lib/CodeGen/CodeGenFunction.h clang/lib/Parse/ParseOpenACC.cpp clang/lib/Sema/SemaExceptionSpec.cpp clang/lib/Sema/SemaOpenACC.cpp clang/lib/Sema/TreeTransform.h clang/lib/Serialization/ASTReaderStmt.cpp clang/lib/Serialization/ASTWriterStmt.cpp clang/lib/StaticAnalyzer/Core/ExprEngine.cpp clang/test/ParserOpenACC/parse-cache-construct.c clang/test/ParserOpenACC/parse-cache-construct.cpp clang/tools/libclang/CIndex.cpp clang/tools/libclang/CXCursor.cpp Removed: ################################################################################ diff --git a/clang/include/clang-c/Index.h b/clang/include/clang-c/Index.h index c50410dc365b6..38e2417dcd181 100644 --- a/clang/include/clang-c/Index.h +++ b/clang/include/clang-c/Index.h @@ -2214,7 +2214,11 @@ enum CXCursorKind { */ CXCursor_OpenACCAtomicConstruct = 332, - CXCursor_LastStmt = CXCursor_OpenACCAtomicConstruct, + /** OpenACC cache Construct. + */ + CXCursor_OpenACCCacheConstruct = 333, + + CXCursor_LastStmt = CXCursor_OpenACCCacheConstruct, /** * Cursor that represents the translation unit itself. diff --git a/clang/include/clang/AST/RecursiveASTVisitor.h b/clang/include/clang/AST/RecursiveASTVisitor.h index b1b4363b65721..89757ddb0781f 100644 --- a/clang/include/clang/AST/RecursiveASTVisitor.h +++ b/clang/include/clang/AST/RecursiveASTVisitor.h @@ -4113,6 +4113,10 @@ DEF_TRAVERSE_STMT(OpenACCUpdateConstruct, { TRY_TO(VisitOpenACCClauseList(S->clauses())); }) DEF_TRAVERSE_STMT(OpenACCAtomicConstruct, { TRY_TO(TraverseOpenACCAssociatedStmtConstruct(S)); }) +DEF_TRAVERSE_STMT(OpenACCCacheConstruct, { + for (auto *E : S->getVarList()) + TRY_TO(TraverseStmt(E)); +}) // Traverse HLSL: Out argument expression DEF_TRAVERSE_STMT(HLSLOutArgExpr, {}) diff --git a/clang/include/clang/AST/StmtOpenACC.h b/clang/include/clang/AST/StmtOpenACC.h index bd6c95d342ce2..c2c74f5cf1958 100644 --- a/clang/include/clang/AST/StmtOpenACC.h +++ b/clang/include/clang/AST/StmtOpenACC.h @@ -593,6 +593,81 @@ class OpenACCWaitConstruct final } }; +class OpenACCCacheConstruct final + : public OpenACCConstructStmt, + private llvm::TrailingObjects<OpenACCCacheConstruct, Expr *> { + friend TrailingObjects; + friend class ASTStmtWriter; + friend class ASTStmtReader; + // Locations of the left and right parens of the 'var-list' + // expression-list. + SourceRange ParensLoc; + SourceLocation ReadOnlyLoc; + + unsigned NumVars = 0; + + OpenACCCacheConstruct(unsigned NumVars) + : OpenACCConstructStmt(OpenACCCacheConstructClass, + OpenACCDirectiveKind::Cache, SourceLocation{}, + SourceLocation{}, SourceLocation{}), + NumVars(NumVars) { + std::uninitialized_value_construct(getVarListPtr(), + getVarListPtr() + NumVars); + } + OpenACCCacheConstruct(SourceLocation Start, SourceLocation DirectiveLoc, + SourceLocation LParenLoc, SourceLocation ReadOnlyLoc, + ArrayRef<Expr *> VarList, SourceLocation RParenLoc, + SourceLocation End) + : OpenACCConstructStmt(OpenACCCacheConstructClass, + OpenACCDirectiveKind::Cache, Start, DirectiveLoc, + End), + ParensLoc(LParenLoc, RParenLoc), ReadOnlyLoc(ReadOnlyLoc), + NumVars(VarList.size()) { + + std::uninitialized_copy(VarList.begin(), VarList.end(), getVarListPtr()); + } + + Expr **getVarListPtr() const { + return const_cast<Expr **>(getTrailingObjects<Expr *>()); + } + +public: + llvm::ArrayRef<Expr *> getVarList() const { + return llvm::ArrayRef<Expr *>(getVarListPtr(), NumVars); + } + + llvm::ArrayRef<Expr *> getVarList() { + return llvm::ArrayRef<Expr *>(getVarListPtr(), NumVars); + } + + static bool classof(const Stmt *T) { + return T->getStmtClass() == OpenACCCacheConstructClass; + } + + static OpenACCCacheConstruct *CreateEmpty(const ASTContext &C, + unsigned NumVars); + static OpenACCCacheConstruct * + Create(const ASTContext &C, SourceLocation Start, SourceLocation DirectiveLoc, + SourceLocation LParenLoc, SourceLocation ReadOnlyLoc, + ArrayRef<Expr *> VarList, SourceLocation RParenLoc, + SourceLocation End); + + SourceLocation getLParenLoc() const { return ParensLoc.getBegin(); } + SourceLocation getRParenLoc() const { return ParensLoc.getEnd(); } + bool hasReadOnly() const { return !ReadOnlyLoc.isInvalid(); } + SourceLocation getReadOnlyLoc() const { return ReadOnlyLoc; } + + child_range children() { + Stmt **Begin = reinterpret_cast<Stmt **>(getVarListPtr()); + return child_range(Begin, Begin + NumVars); + } + + const_child_range children() const { + Stmt *const *Begin = reinterpret_cast<Stmt *const *>(getVarListPtr()); + return const_child_range(Begin, Begin + NumVars); + } +}; + // This class represents an 'init' construct, which has just a clause list. class OpenACCInitConstruct final : public OpenACCConstructStmt, diff --git a/clang/include/clang/AST/TextNodeDumper.h b/clang/include/clang/AST/TextNodeDumper.h index b0e5b28e22863..0925aca783bba 100644 --- a/clang/include/clang/AST/TextNodeDumper.h +++ b/clang/include/clang/AST/TextNodeDumper.h @@ -422,6 +422,7 @@ class TextNodeDumper void VisitOpenACCShutdownConstruct(const OpenACCShutdownConstruct *S); void VisitOpenACCUpdateConstruct(const OpenACCUpdateConstruct *S); void VisitOpenACCAtomicConstruct(const OpenACCAtomicConstruct *S); + void VisitOpenACCCacheConstruct(const OpenACCCacheConstruct *S); void VisitOpenACCAsteriskSizeExpr(const OpenACCAsteriskSizeExpr *S); void VisitOpenACCDeclareDecl(const OpenACCDeclareDecl *D); void VisitEmbedExpr(const EmbedExpr *S); diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 0efb15405ed5d..d89648a8a2e83 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12848,6 +12848,9 @@ def err_acc_not_a_var_ref def err_acc_not_a_var_ref_use_device_declare : Error<"OpenACC variable %select{in 'use_device' clause|on 'declare' " "construct}0 is not a valid variable name or array name">; +def err_acc_not_a_var_ref_cache + : Error<"OpenACC variable in cache directive is not a valid sub-array or " + "array element">; def err_acc_typecheck_subarray_value : Error<"OpenACC sub-array subscripted value is not an array or pointer">; def err_acc_subarray_function_type diff --git a/clang/include/clang/Basic/StmtNodes.td b/clang/include/clang/Basic/StmtNodes.td index ae49671058a01..9526fa5808aa5 100644 --- a/clang/include/clang/Basic/StmtNodes.td +++ b/clang/include/clang/Basic/StmtNodes.td @@ -320,6 +320,7 @@ def OpenACCShutdownConstruct : StmtNode<OpenACCConstructStmt>; def OpenACCSetConstruct : StmtNode<OpenACCConstructStmt>; def OpenACCUpdateConstruct : StmtNode<OpenACCConstructStmt>; def OpenACCAtomicConstruct : StmtNode<OpenACCAssociatedStmtConstruct>; +def OpenACCCacheConstruct : StmtNode<OpenACCConstructStmt>; // OpenACC Additional Expressions. def OpenACCAsteriskSizeExpr : StmtNode<Expr>; diff --git a/clang/include/clang/Parse/Parser.h b/clang/include/clang/Parse/Parser.h index 0602f44333b20..049156e266c70 100644 --- a/clang/include/clang/Parse/Parser.h +++ b/clang/include/clang/Parse/Parser.h @@ -3730,6 +3730,11 @@ class Parser : public CodeCompletionHandler { return Out; } }; + struct OpenACCCacheParseInfo { + bool Failed = false; + SourceLocation ReadOnlyLoc; + SmallVector<Expr *> Vars; + }; /// Represents the 'error' state of parsing an OpenACC Clause, and stores /// whether we can continue parsing, or should give up on the directive. @@ -3752,7 +3757,7 @@ class Parser : public CodeCompletionHandler { /// Helper that parses an ID Expression based on the language options. ExprResult ParseOpenACCIDExpression(); /// Parses the variable list for the `cache` construct. - void ParseOpenACCCacheVarList(); + OpenACCCacheParseInfo ParseOpenACCCacheVarList(); using OpenACCVarParseResult = std::pair<ExprResult, OpenACCParseCanContinue>; /// Parses a single variable in a variable list for OpenACC. diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h index 36a2b6ff7edc3..6edc0d6dd5653 100644 --- a/clang/include/clang/Sema/SemaOpenACC.h +++ b/clang/include/clang/Sema/SemaOpenACC.h @@ -762,6 +762,8 @@ class SemaOpenACC : public SemaBase { /// declaration reference to a variable of the correct type. ExprResult ActOnVar(OpenACCDirectiveKind DK, OpenACCClauseKind CK, Expr *VarExpr); + /// Helper function called by ActonVar that is used to check a 'cache' var. + ExprResult ActOnCacheVar(Expr *VarExpr); // Called after 'ActOnVar' specifically for a 'link' clause, which has to do // some minor additional checks. diff --git a/clang/include/clang/Serialization/ASTBitCodes.h b/clang/include/clang/Serialization/ASTBitCodes.h index 5698bebf13445..efb52cfd536af 100644 --- a/clang/include/clang/Serialization/ASTBitCodes.h +++ b/clang/include/clang/Serialization/ASTBitCodes.h @@ -2049,6 +2049,7 @@ enum StmtCode { STMT_OPENACC_SET_CONSTRUCT, STMT_OPENACC_UPDATE_CONSTRUCT, STMT_OPENACC_ATOMIC_CONSTRUCT, + STMT_OPENACC_CACHE_CONSTRUCT, // HLSL Constructs EXPR_HLSL_OUT_ARG, diff --git a/clang/lib/AST/StmtOpenACC.cpp b/clang/lib/AST/StmtOpenACC.cpp index 11eab0c27579d..8a86074fe68a0 100644 --- a/clang/lib/AST/StmtOpenACC.cpp +++ b/clang/lib/AST/StmtOpenACC.cpp @@ -321,3 +321,21 @@ OpenACCAtomicConstruct *OpenACCAtomicConstruct::Create( OpenACCAtomicConstruct(Start, DirectiveLoc, AtKind, End, AssociatedStmt); return Inst; } +OpenACCCacheConstruct *OpenACCCacheConstruct::CreateEmpty(const ASTContext &C, + unsigned NumVars) { + void *Mem = + C.Allocate(OpenACCCacheConstruct::totalSizeToAlloc<Expr *>(NumVars)); + auto *Inst = new (Mem) OpenACCCacheConstruct(NumVars); + return Inst; +} + +OpenACCCacheConstruct *OpenACCCacheConstruct::Create( + const ASTContext &C, SourceLocation Start, SourceLocation DirectiveLoc, + SourceLocation LParenLoc, SourceLocation ReadOnlyLoc, + ArrayRef<Expr *> VarList, SourceLocation RParenLoc, SourceLocation End) { + void *Mem = C.Allocate( + OpenACCCacheConstruct::totalSizeToAlloc<Expr *>(VarList.size())); + auto *Inst = new (Mem) OpenACCCacheConstruct( + Start, DirectiveLoc, LParenLoc, ReadOnlyLoc, VarList, RParenLoc, End); + return Inst; +} diff --git a/clang/lib/AST/StmtPrinter.cpp b/clang/lib/AST/StmtPrinter.cpp index dac9a95e10f3d..263c1bae8e075 100644 --- a/clang/lib/AST/StmtPrinter.cpp +++ b/clang/lib/AST/StmtPrinter.cpp @@ -1261,6 +1261,18 @@ void StmtPrinter::VisitOpenACCAtomicConstruct(OpenACCAtomicConstruct *S) { PrintStmt(S->getAssociatedStmt()); } +void StmtPrinter::VisitOpenACCCacheConstruct(OpenACCCacheConstruct *S) { + Indent() << "#pragma acc cache("; + if (S->hasReadOnly()) + OS << "readonly: "; + + llvm::interleaveComma(S->getVarList(), OS, [&](const Expr *E) { + E->printPretty(OS, nullptr, Policy); + }); + + OS << ")\n"; +} + //===----------------------------------------------------------------------===// // Expr printing methods. //===----------------------------------------------------------------------===// diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp index e283a9ad4a567..f9aa7aa079e67 100644 --- a/clang/lib/AST/StmtProfile.cpp +++ b/clang/lib/AST/StmtProfile.cpp @@ -2708,6 +2708,7 @@ void OpenACCClauseProfiler::VisitWaitClause(const OpenACCWaitClause &Clause) { for (auto *E : Clause.getQueueIdExprs()) Profiler.VisitStmt(E); } + /// Nothing to do here, there are no sub-statements. void OpenACCClauseProfiler::VisitDeviceTypeClause( const OpenACCDeviceTypeClause &Clause) {} @@ -2796,6 +2797,11 @@ void StmtProfiler::VisitOpenACCWaitConstruct(const OpenACCWaitConstruct *S) { P.VisitOpenACCClauseList(S->clauses()); } +void StmtProfiler::VisitOpenACCCacheConstruct(const OpenACCCacheConstruct *S) { + // VisitStmt covers 'children', so the exprs inside of it are covered. + VisitStmt(S); +} + void StmtProfiler::VisitOpenACCInitConstruct(const OpenACCInitConstruct *S) { VisitStmt(S); OpenACCClauseProfiler P{*this}; diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp index 063e718454d46..9b2abce5421a6 100644 --- a/clang/lib/AST/TextNodeDumper.cpp +++ b/clang/lib/AST/TextNodeDumper.cpp @@ -3042,6 +3042,12 @@ void TextNodeDumper::VisitOpenACCHostDataConstruct( void TextNodeDumper::VisitOpenACCWaitConstruct(const OpenACCWaitConstruct *S) { VisitOpenACCConstructStmt(S); } +void TextNodeDumper::VisitOpenACCCacheConstruct( + const OpenACCCacheConstruct *S) { + VisitOpenACCConstructStmt(S); + if (S->hasReadOnly()) + OS <<" readonly"; +} void TextNodeDumper::VisitOpenACCInitConstruct(const OpenACCInitConstruct *S) { VisitOpenACCConstructStmt(S); } diff --git a/clang/lib/CodeGen/CGStmt.cpp b/clang/lib/CodeGen/CGStmt.cpp index e56ba6c3e8803..abe799af32c6e 100644 --- a/clang/lib/CodeGen/CGStmt.cpp +++ b/clang/lib/CodeGen/CGStmt.cpp @@ -494,6 +494,10 @@ void CodeGenFunction::EmitStmt(const Stmt *S, ArrayRef<const Attr *> Attrs) { break; case Stmt::OpenACCAtomicConstructClass: EmitOpenACCAtomicConstruct(cast<OpenACCAtomicConstruct>(*S)); + break; + case Stmt::OpenACCCacheConstructClass: + EmitOpenACCCacheConstruct(cast<OpenACCCacheConstruct>(*S)); + break; } } diff --git a/clang/lib/CodeGen/CodeGenFunction.h b/clang/lib/CodeGen/CodeGenFunction.h index 7c0d6c3685597..018fc66b72a1e 100644 --- a/clang/lib/CodeGen/CodeGenFunction.h +++ b/clang/lib/CodeGen/CodeGenFunction.h @@ -4192,6 +4192,10 @@ class CodeGenFunction : public CodeGenTypeCache { // some sort of IR. EmitStmt(S.getAssociatedStmt()); } + void EmitOpenACCCacheConstruct(const OpenACCCacheConstruct &S) { + // TODO OpenACC: Implement this. It is currently implemented as a 'no-op', + // but in the future we will implement some sort of IR. + } //===--------------------------------------------------------------------===// // LValue Expression Emission diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp index 5eefd3ffb092e..1417caae0bbde 100644 --- a/clang/lib/Parse/ParseOpenACC.cpp +++ b/clang/lib/Parse/ParseOpenACC.cpp @@ -617,13 +617,18 @@ unsigned getOpenACCScopeFlags(OpenACCDirectiveKind DirKind) { case OpenACCDirectiveKind::Wait: case OpenACCDirectiveKind::Init: case OpenACCDirectiveKind::Shutdown: + case OpenACCDirectiveKind::Cache: + case OpenACCDirectiveKind::Loop: + case OpenACCDirectiveKind::Atomic: + case OpenACCDirectiveKind::Declare: + case OpenACCDirectiveKind::Routine: + case OpenACCDirectiveKind::Set: + case OpenACCDirectiveKind::Update: return 0; case OpenACCDirectiveKind::Invalid: llvm_unreachable("Shouldn't be creating a scope for an invalid construct"); - default: - break; } - return 0; + llvm_unreachable("Shouldn't be creating a scope for an invalid construct"); } } // namespace @@ -1403,25 +1408,29 @@ llvm::SmallVector<Expr *> Parser::ParseOpenACCVarList(OpenACCDirectiveKind DK, /// In C and C++, the syntax of the cache directive is: /// /// #pragma acc cache ([readonly:]var-list) new-line -void Parser::ParseOpenACCCacheVarList() { +Parser::OpenACCCacheParseInfo Parser::ParseOpenACCCacheVarList() { // If this is the end of the line, just return 'false' and count on the close // paren diagnostic to catch the issue. if (getCurToken().isAnnotation()) - return; + return {}; + + OpenACCCacheParseInfo CacheInfo; + SourceLocation ReadOnlyLoc = getCurToken().getLocation(); // The VarList is an optional `readonly:` followed by a list of a variable // specifications. Consume something that looks like a 'tag', and diagnose if // it isn't 'readonly'. if (tryParseAndConsumeSpecialTokenKind(*this, OpenACCSpecialTokenKind::ReadOnly, - OpenACCDirectiveKind::Cache)) { - // FIXME: Record that this is a 'readonly' so that we can use that during - // Sema/AST generation. - } + OpenACCDirectiveKind::Cache)) + CacheInfo.ReadOnlyLoc = ReadOnlyLoc; // ParseOpenACCVarList should leave us before a r-paren, so no need to skip // anything here. - ParseOpenACCVarList(OpenACCDirectiveKind::Cache, OpenACCClauseKind::Invalid); + CacheInfo.Vars = ParseOpenACCVarList(OpenACCDirectiveKind::Cache, + OpenACCClauseKind::Invalid); + + return CacheInfo; } Parser::OpenACCDirectiveParseInfo @@ -1430,6 +1439,7 @@ Parser::ParseOpenACCDirective() { SourceLocation DirLoc = getCurToken().getLocation(); OpenACCDirectiveKind DirKind = ParseOpenACCDirectiveKind(*this); Parser::OpenACCWaitParseInfo WaitInfo; + Parser::OpenACCCacheParseInfo CacheInfo; OpenACCAtomicKind AtomicKind = OpenACCAtomicKind::None; getActions().OpenACC().ActOnConstruct(DirKind, DirLoc); @@ -1464,7 +1474,7 @@ Parser::ParseOpenACCDirective() { break; } case OpenACCDirectiveKind::Cache: - ParseOpenACCCacheVarList(); + CacheInfo = ParseOpenACCCacheVarList(); // The ParseOpenACCCacheVarList function manages to recover from failures, // so we can always consume the close. T.consumeClose(); @@ -1486,16 +1496,19 @@ Parser::ParseOpenACCDirective() { } // Parses the list of clauses, if present, plus set up return value. - OpenACCDirectiveParseInfo ParseInfo{DirKind, - StartLoc, - DirLoc, - T.getOpenLocation(), - T.getCloseLocation(), - /*EndLoc=*/SourceLocation{}, - WaitInfo.QueuesLoc, - AtomicKind, - WaitInfo.getAllExprs(), - ParseOpenACCClauseList(DirKind)}; + OpenACCDirectiveParseInfo ParseInfo{ + DirKind, + StartLoc, + DirLoc, + T.getOpenLocation(), + T.getCloseLocation(), + /*EndLoc=*/SourceLocation{}, + (DirKind == OpenACCDirectiveKind::Wait ? WaitInfo.QueuesLoc + : CacheInfo.ReadOnlyLoc), + AtomicKind, + (DirKind == OpenACCDirectiveKind::Wait ? WaitInfo.getAllExprs() + : CacheInfo.Vars), + ParseOpenACCClauseList(DirKind)}; assert(Tok.is(tok::annot_pragma_openacc_end) && "Didn't parse all OpenACC Clauses"); diff --git a/clang/lib/Sema/SemaExceptionSpec.cpp b/clang/lib/Sema/SemaExceptionSpec.cpp index a8eb24133a76d..f358d8342e2f3 100644 --- a/clang/lib/Sema/SemaExceptionSpec.cpp +++ b/clang/lib/Sema/SemaExceptionSpec.cpp @@ -1407,6 +1407,7 @@ CanThrowResult Sema::canThrow(const Stmt *S) { case Stmt::OpenACCEnterDataConstructClass: case Stmt::OpenACCExitDataConstructClass: case Stmt::OpenACCWaitConstructClass: + case Stmt::OpenACCCacheConstructClass: case Stmt::OpenACCInitConstructClass: case Stmt::OpenACCShutdownConstructClass: case Stmt::OpenACCSetConstructClass: diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp index 59583e73952cf..85fa631d61e89 100644 --- a/clang/lib/Sema/SemaOpenACC.cpp +++ b/clang/lib/Sema/SemaOpenACC.cpp @@ -81,6 +81,9 @@ bool PreserveLoopRAIIDepthInAssociatedStmtRAII(OpenACCDirectiveKind DK) { case OpenACCDirectiveKind::HostData: case OpenACCDirectiveKind::Atomic: return true; + case OpenACCDirectiveKind::Cache: + case OpenACCDirectiveKind::Routine: + case OpenACCDirectiveKind::Declare: case OpenACCDirectiveKind::EnterData: case OpenACCDirectiveKind::ExitData: case OpenACCDirectiveKind::Wait: @@ -89,7 +92,6 @@ bool PreserveLoopRAIIDepthInAssociatedStmtRAII(OpenACCDirectiveKind DK) { case OpenACCDirectiveKind::Set: case OpenACCDirectiveKind::Update: llvm_unreachable("Doesn't have an associated stmt"); - default: case OpenACCDirectiveKind::Invalid: llvm_unreachable("Unhandled directive kind?"); } @@ -334,6 +336,7 @@ void SemaOpenACC::ActOnConstruct(OpenACCDirectiveKind K, case OpenACCDirectiveKind::Shutdown: case OpenACCDirectiveKind::Set: case OpenACCDirectiveKind::Update: + case OpenACCDirectiveKind::Cache: case OpenACCDirectiveKind::Atomic: case OpenACCDirectiveKind::Declare: // Nothing to do here, there is no real legalization that needs to happen @@ -480,8 +483,56 @@ bool SemaOpenACC::CheckVarIsPointerType(OpenACCClauseKind ClauseKind, return false; } +ExprResult SemaOpenACC::ActOnCacheVar(Expr *VarExpr) { + Expr *CurVarExpr = VarExpr->IgnoreParenImpCasts(); + if (!isa<ArraySectionExpr, ArraySubscriptExpr>(CurVarExpr)) { + Diag(VarExpr->getExprLoc(), diag::err_acc_not_a_var_ref_cache); + return ExprError(); + } + + // It isn't clear what 'simple array element or simple subarray' means, so we + // will just allow arbitrary depth. + while (isa<ArraySectionExpr, ArraySubscriptExpr>(CurVarExpr)) { + if (auto *SubScrpt = dyn_cast<ArraySubscriptExpr>(CurVarExpr)) + CurVarExpr = SubScrpt->getBase()->IgnoreParenImpCasts(); + else + CurVarExpr = + cast<ArraySectionExpr>(CurVarExpr)->getBase()->IgnoreParenImpCasts(); + } + + // References to a VarDecl are fine. + if (const auto *DRE = dyn_cast<DeclRefExpr>(CurVarExpr)) { + if (isa<VarDecl, NonTypeTemplateParmDecl>( + DRE->getFoundDecl()->getCanonicalDecl())) + return VarExpr; + } + + if (const auto *ME = dyn_cast<MemberExpr>(CurVarExpr)) { + if (isa<FieldDecl>(ME->getMemberDecl()->getCanonicalDecl())) { + return VarExpr; + } + } + + // Nothing really we can do here, as these are dependent. So just return they + // are valid. + if (isa<DependentScopeDeclRefExpr, CXXDependentScopeMemberExpr>(CurVarExpr)) + return VarExpr; + + // There isn't really anything we can do in the case of a recovery expr, so + // skip the diagnostic rather than produce a confusing diagnostic. + if (isa<RecoveryExpr>(CurVarExpr)) + return ExprError(); + + Diag(VarExpr->getExprLoc(), diag::err_acc_not_a_var_ref_cache); + return ExprError(); +} ExprResult SemaOpenACC::ActOnVar(OpenACCDirectiveKind DK, OpenACCClauseKind CK, Expr *VarExpr) { + // This has unique enough restrictions that we should split it to a separate + // function. + if (DK == OpenACCDirectiveKind::Cache) + return ActOnCacheVar(VarExpr); + Expr *CurVarExpr = VarExpr->IgnoreParenImpCasts(); // 'use_device' doesn't allow array subscript or array sections. @@ -1624,6 +1675,12 @@ StmtResult SemaOpenACC::ActOnEndStmtDirective( getASTContext(), StartLoc, DirLoc, AtomicKind, EndLoc, AssocStmt.isUsable() ? AssocStmt.get() : nullptr); } + case OpenACCDirectiveKind::Cache: { + assert(Clauses.empty() && "Cache doesn't allow clauses"); + return OpenACCCacheConstruct::Create(getASTContext(), StartLoc, DirLoc, + LParenLoc, MiscLoc, Exprs, RParenLoc, + EndLoc); + } case OpenACCDirectiveKind::Declare: { // Declare is a declaration directive, but can be used here as long as we // wrap it in a DeclStmt. So make sure we do that here. @@ -1649,6 +1706,7 @@ StmtResult SemaOpenACC::ActOnAssociatedStmt( case OpenACCDirectiveKind::Init: case OpenACCDirectiveKind::Shutdown: case OpenACCDirectiveKind::Set: + case OpenACCDirectiveKind::Cache: llvm_unreachable( "these don't have associated statements, so shouldn't get here"); case OpenACCDirectiveKind::Atomic: diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h index ba56179cad7b1..cb2f335e4290f 100644 --- a/clang/lib/Sema/TreeTransform.h +++ b/clang/lib/Sema/TreeTransform.h @@ -4204,6 +4204,15 @@ class TreeTransform { Exprs, RParenLoc, EndLoc, Clauses, {}); } + StmtResult RebuildOpenACCCacheConstruct( + SourceLocation BeginLoc, SourceLocation DirLoc, SourceLocation LParenLoc, + SourceLocation ReadOnlyLoc, ArrayRef<Expr *> VarList, + SourceLocation RParenLoc, SourceLocation EndLoc) { + return getSema().OpenACC().ActOnEndStmtDirective( + OpenACCDirectiveKind::Cache, BeginLoc, DirLoc, LParenLoc, ReadOnlyLoc, + VarList, RParenLoc, EndLoc, {}, {}); + } + StmtResult RebuildOpenACCAtomicConstruct(SourceLocation BeginLoc, SourceLocation DirLoc, OpenACCAtomicKind AtKind, @@ -12646,6 +12655,37 @@ TreeTransform<Derived>::TransformOpenACCWaitConstruct(OpenACCWaitConstruct *C) { DevNumExpr.isUsable() ? DevNumExpr.get() : nullptr, C->getQueuesLoc(), QueueIdExprs, C->getRParenLoc(), C->getEndLoc(), TransformedClauses); } +template <typename Derived> +StmtResult TreeTransform<Derived>::TransformOpenACCCacheConstruct( + OpenACCCacheConstruct *C) { + getSema().OpenACC().ActOnConstruct(C->getDirectiveKind(), C->getBeginLoc()); + + llvm::SmallVector<Expr *> TransformedVarList; + for (Expr *Var : C->getVarList()) { + assert(Var && "Null var listexpr?"); + + ExprResult NewVar = getDerived().TransformExpr(Var); + + if (!NewVar.isUsable()) + break; + + NewVar = getSema().OpenACC().ActOnVar( + C->getDirectiveKind(), OpenACCClauseKind::Invalid, NewVar.get()); + if (!NewVar.isUsable()) + break; + + TransformedVarList.push_back(NewVar.get()); + } + + if (getSema().OpenACC().ActOnStartStmtDirective(C->getDirectiveKind(), + C->getBeginLoc(), {})) + return StmtError(); + + return getDerived().RebuildOpenACCCacheConstruct( + C->getBeginLoc(), C->getDirectiveLoc(), C->getLParenLoc(), + C->getReadOnlyLoc(), TransformedVarList, C->getRParenLoc(), + C->getEndLoc()); +} template <typename Derived> StmtResult TreeTransform<Derived>::TransformOpenACCAtomicConstruct( diff --git a/clang/lib/Serialization/ASTReaderStmt.cpp b/clang/lib/Serialization/ASTReaderStmt.cpp index 835ad4a658944..48f9f89bd6e4c 100644 --- a/clang/lib/Serialization/ASTReaderStmt.cpp +++ b/clang/lib/Serialization/ASTReaderStmt.cpp @@ -2922,6 +2922,16 @@ void ASTStmtReader::VisitOpenACCWaitConstruct(OpenACCWaitConstruct *S) { } } +void ASTStmtReader::VisitOpenACCCacheConstruct(OpenACCCacheConstruct *S) { + VisitStmt(S); + (void)Record.readInt(); + VisitOpenACCConstructStmt(S); + S->ParensLoc = Record.readSourceRange(); + S->ReadOnlyLoc = Record.readSourceLocation(); + for (unsigned I = 0; I < S->NumVars; ++I) + S->getVarListPtr()[I] = cast<Expr>(Record.readSubStmt()); +} + void ASTStmtReader::VisitOpenACCAtomicConstruct(OpenACCAtomicConstruct *S) { VisitStmt(S); S->Kind = Record.readEnum<OpenACCDirectiveKind>(); @@ -4447,6 +4457,11 @@ Stmt *ASTReader::ReadStmtFromStream(ModuleFile &F) { S = OpenACCWaitConstruct::CreateEmpty(Context, NumExprs, NumClauses); break; } + case STMT_OPENACC_CACHE_CONSTRUCT: { + unsigned NumVars = Record[ASTStmtReader::NumStmtFields]; + S = OpenACCCacheConstruct::CreateEmpty(Context, NumVars); + break; + } case STMT_OPENACC_INIT_CONSTRUCT: { unsigned NumClauses = Record[ASTStmtReader::NumStmtFields]; S = OpenACCInitConstruct::CreateEmpty(Context, NumClauses); diff --git a/clang/lib/Serialization/ASTWriterStmt.cpp b/clang/lib/Serialization/ASTWriterStmt.cpp index 82738d3a8c88a..aa5a7854394a0 100644 --- a/clang/lib/Serialization/ASTWriterStmt.cpp +++ b/clang/lib/Serialization/ASTWriterStmt.cpp @@ -3017,6 +3017,18 @@ void ASTStmtWriter::VisitOpenACCAtomicConstruct(OpenACCAtomicConstruct *S) { Code = serialization::STMT_OPENACC_ATOMIC_CONSTRUCT; } +void ASTStmtWriter::VisitOpenACCCacheConstruct(OpenACCCacheConstruct *S) { + VisitStmt(S); + Record.push_back(S->getVarList().size()); + VisitOpenACCConstructStmt(S); + Record.AddSourceRange(S->ParensLoc); + Record.AddSourceLocation(S->ReadOnlyLoc); + + for (Expr *E : S->getVarList()) + Record.AddStmt(E); + Code = serialization::STMT_OPENACC_CACHE_CONSTRUCT; +} + //===----------------------------------------------------------------------===// // HLSL Constructs/Directives. //===----------------------------------------------------------------------===// diff --git a/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp b/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp index 318fa3c1caf06..b76bbf72768ae 100644 --- a/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp +++ b/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp @@ -1834,6 +1834,7 @@ void ExprEngine::Visit(const Stmt *S, ExplodedNode *Pred, case Stmt::OpenACCExitDataConstructClass: case Stmt::OpenACCHostDataConstructClass: case Stmt::OpenACCWaitConstructClass: + case Stmt::OpenACCCacheConstructClass: case Stmt::OpenACCInitConstructClass: case Stmt::OpenACCShutdownConstructClass: case Stmt::OpenACCSetConstructClass: diff --git a/clang/test/AST/ast-print-openacc-cache-construct.cpp b/clang/test/AST/ast-print-openacc-cache-construct.cpp new file mode 100644 index 0000000000000..e98ff3ca67600 --- /dev/null +++ b/clang/test/AST/ast-print-openacc-cache-construct.cpp @@ -0,0 +1,9 @@ +// RUN: %clang_cc1 -fopenacc -ast-print %s -o - | FileCheck %s + +void foo() { + int Array[5]; + // CHECK: #pragma acc cache(readonly: Array[1], Array[1:2]) + #pragma acc cache(readonly:Array[1], Array[1:2]) + // CHECK: #pragma acc cache(Array[1], Array[1:2]) + #pragma acc cache(Array[1], Array[1:2]) +} diff --git a/clang/test/ParserOpenACC/parse-cache-construct.c b/clang/test/ParserOpenACC/parse-cache-construct.c index 8937aa095d5ea..d52a5e4ea3441 100644 --- a/clang/test/ParserOpenACC/parse-cache-construct.c +++ b/clang/test/ParserOpenACC/parse-cache-construct.c @@ -12,186 +12,165 @@ void func() { struct S s; for (int i = 0; i < 10; ++i) { - // expected-error@+2{{expected '('}} - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+1{{expected '('}} #pragma acc cache } for (int i = 0; i < 10; ++i) { - // expected-error@+3{{expected '('}} - // expected-error@+2{{invalid OpenACC clause 'clause'}} - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+2{{expected '('}} + // expected-error@+1{{invalid OpenACC clause 'clause'}} #pragma acc cache clause list } for (int i = 0; i < 10; ++i) { - // expected-error@+2{{expected expression}} - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+1{{expected expression}} #pragma acc cache() } for (int i = 0; i < 10; ++i) { - // expected-error@+3{{expected expression}} - // expected-error@+2{{invalid OpenACC clause 'clause'}} - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+2{{expected expression}} + // expected-error@+1{{invalid OpenACC clause 'clause'}} #pragma acc cache() clause-list } for (int i = 0; i < 10; ++i) { - // expected-error@+3{{expected ')'}} - // expected-note@+2{{to match this '('}} - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+2{{expected ')'}} + // expected-note@+1{{to match this '('}} #pragma acc cache( } for (int i = 0; i < 10; ++i) { - // expected-error@+4{{use of undeclared identifier 'invalid'}} - // expected-error@+3{{expected ')'}} - // expected-note@+2{{to match this '('}} - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+3{{use of undeclared identifier 'invalid'}} + // expected-error@+2{{expected ')'}} + // expected-note@+1{{to match this '('}} #pragma acc cache(invalid } for (int i = 0; i < 10; ++i) { // expected-error@+3{{expected ')'}} // expected-note@+2{{to match this '('}} - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+1{{OpenACC variable in cache directive is not a valid sub-array or array element}} #pragma acc cache(ArrayPtr } for (int i = 0; i < 10; ++i) { - // expected-error@+2{{use of undeclared identifier 'invalid'}} - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+1{{use of undeclared identifier 'invalid'}} #pragma acc cache(invalid) } for (int i = 0; i < 10; ++i) { - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+1{{OpenACC variable in cache directive is not a valid sub-array or array element}} #pragma acc cache(ArrayPtr) } for (int i = 0; i < 10; ++i) { - // expected-error@+6{{expected expression}} - // expected-error@+5{{expected ']'}} - // expected-note@+4{{to match this '['}} - // expected-error@+3{{expected ')'}} - // expected-note@+2{{to match this '('}} - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+5{{expected expression}} + // expected-error@+4{{expected ']'}} + // expected-note@+3{{to match this '['}} + // expected-error@+2{{expected ')'}} + // expected-note@+1{{to match this '('}} #pragma acc cache(ArrayPtr[ } for (int i = 0; i < 10; ++i) { - // expected-error@+4{{expected expression}} - // expected-error@+3{{expected ']'}} - // expected-note@+2{{to match this '['}} - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+3{{expected expression}} + // expected-error@+2{{expected ']'}} + // expected-note@+1{{to match this '['}} #pragma acc cache(ArrayPtr[, 5) } for (int i = 0; i < 10; ++i) { - // expected-error@+4{{expected expression}} - // expected-error@+3{{expected ']'}} - // expected-note@+2{{to match this '['}} - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+3{{expected expression}} + // expected-error@+2{{expected ']'}} + // expected-note@+1{{to match this '['}} #pragma acc cache(Array[) } for (int i = 0; i < 10; ++i) { - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} #pragma acc cache(Array[*readonly]) } for (int i = 0; i < 10; ++i) { - // expected-error@+6{{expected expression}} - // expected-error@+5{{expected ']'}} - // expected-note@+4{{to match this '['}} - // expected-error@+3{{expected ')'}} - // expected-note@+2{{to match this '('}} - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+5{{expected expression}} + // expected-error@+4{{expected ']'}} + // expected-note@+3{{to match this '['}} + // expected-error@+2{{expected ')'}} + // expected-note@+1{{to match this '('}} #pragma acc cache(Array[*readonly: } for (int i = 0; i < 10; ++i) { - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+1{{OpenACC variable in cache directive is not a valid sub-array or array element}} #pragma acc cache(readonly) } for (int i = 0; i < 10; ++i) { // expected-error@+2{{invalid tag 'devnum' on 'cache' directive}} - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+1{{OpenACC variable in cache directive is not a valid sub-array or array element}} #pragma acc cache(devnum:ArrayPtr) } for (int i = 0; i < 10; ++i) { // expected-error@+2{{invalid tag 'invalid' on 'cache' directive}} - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+1{{OpenACC variable in cache directive is not a valid sub-array or array element}} #pragma acc cache(invalid:ArrayPtr) } for (int i = 0; i < 10; ++i) { - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+1{{OpenACC variable in cache directive is not a valid sub-array or array element}} #pragma acc cache(readonly:ArrayPtr) } for (int i = 0; i < 10; ++i) { - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} #pragma acc cache(readonly:ArrayPtr[5:1]) } for (int i = 0; i < 10; ++i) { - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} #pragma acc cache(readonly:ArrayPtr[5:*readonly]) } for (int i = 0; i < 10; ++i) { - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+1{{OpenACC variable in cache directive is not a valid sub-array or array element}} #pragma acc cache(readonly:ArrayPtr[5:*readonly], Array) } for (int i = 0; i < 10; ++i) { - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} #pragma acc cache(readonly:ArrayPtr[5:*readonly], Array[*readonly:3]) } for (int i = 0; i < 10; ++i) { - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} #pragma acc cache(readonly:ArrayPtr[5 + i:*readonly], Array[*readonly + i:3]) } for (int i = 0; i < 10; ++i) { - // expected-error@+4{{expected expression}} - // expected-error@+3{{expected ')'}} - // expected-note@+2{{to match this '('}} - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+3{{expected expression}} + // expected-error@+2{{expected ')'}} + // expected-note@+1{{to match this '('}} #pragma acc cache(readonly:ArrayPtr[5:*readonly], } for (int i = 0; i < 10; ++i) { - // expected-error@+2{{expected expression}} - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+1{{expected expression}} #pragma acc cache(readonly:ArrayPtr[5:*readonly],) } for (int i = 0; i < 10; ++i) { - // expected-warning@+2{{left operand of comma operator has no effect}} - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-warning@+1{{left operand of comma operator has no effect}} #pragma acc cache(readonly:ArrayPtr[5,6:*readonly]) } for (int i = 0; i < 10; ++i) { - // expected-warning@+2{{left operand of comma operator has no effect}} - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-warning@+1{{left operand of comma operator has no effect}} #pragma acc cache(readonly:ArrayPtr[5:3, *readonly], ArrayPtr[0]) } for (int i = 0; i < 10; ++i) { - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+1{{OpenACC variable in cache directive is not a valid sub-array or array element}} #pragma acc cache(readonly:s.foo) } for (int i = 0; i < 10; ++i) { - // expected-warning@+2{{left operand of comma operator has no effect}} - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-warning@+1{{left operand of comma operator has no effect}} #pragma acc cache(readonly:s.Array[1,2]) } } diff --git a/clang/test/ParserOpenACC/parse-cache-construct.cpp b/clang/test/ParserOpenACC/parse-cache-construct.cpp index 374fe2697b63f..a5a1e58028c33 100644 --- a/clang/test/ParserOpenACC/parse-cache-construct.cpp +++ b/clang/test/ParserOpenACC/parse-cache-construct.cpp @@ -9,34 +9,29 @@ template<typename T, int I> void func() { char *ArrayPtr = getArrayPtr(); for (int i = 0; i < 10; ++i) { - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} - #pragma acc cache(ArrayPtr[T::value + I:I + 5], T::array[(i + T::value, 5): 6]) + // expected-warning@+1{{left operand of comma operator has no effect}} + #pragma acc cache(ArrayPtr[T::value + I:I + 3], T::array[(i + T::value, 2): 4]) } for (int i = 0; i < 10; ++i) { - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} #pragma acc cache(NS::NSArray[NS::NSInt]) } for (int i = 0; i < 10; ++i) { - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} #pragma acc cache(NS::NSArray[NS::NSInt : NS::NSInt]) } for (int i = 0; i < 10; ++i) { - // expected-error@+2{{use of undeclared identifier 'NSArray'; did you mean 'NS::NSArray'}} - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+1{{use of undeclared identifier 'NSArray'; did you mean 'NS::NSArray'}} #pragma acc cache(NSArray[NS::NSInt : NS::NSInt]) } for (int i = 0; i < 10; ++i) { - // expected-error@+2{{use of undeclared identifier 'NSInt'; did you mean 'NS::NSInt'}} - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+1{{use of undeclared identifier 'NSInt'; did you mean 'NS::NSInt'}} #pragma acc cache(NS::NSArray[NSInt : NS::NSInt]) } for (int i = 0; i < 10; ++i) { - // expected-error@+2{{use of undeclared identifier 'NSInt'; did you mean 'NS::NSInt'}} - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+1{{use of undeclared identifier 'NSInt'; did you mean 'NS::NSInt'}} #pragma acc cache(NS::NSArray[NS::NSInt : NSInt]) } } @@ -59,56 +54,46 @@ void use() { Members s; for (int i = 0; i < 10; ++i) { - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} #pragma acc cache(s.array[s.value]) } HasMembersArray Arrs; for (int i = 0; i < 10; ++i) { - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} #pragma acc cache(Arrs.MemArr[3].array[4]) } for (int i = 0; i < 10; ++i) { - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} #pragma acc cache(Arrs.MemArr[3].array[1:4]) } for (int i = 0; i < 10; ++i) { - // expected-error@+2{{OpenACC sub-array is not allowed here}} - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+1{{OpenACC sub-array is not allowed here}} #pragma acc cache(Arrs.MemArr[2:1].array[1:4]) } for (int i = 0; i < 10; ++i) { - // expected-error@+2{{OpenACC sub-array is not allowed here}} - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+1{{OpenACC sub-array is not allowed here}} #pragma acc cache(Arrs.MemArr[2:1].array[4]) } for (int i = 0; i < 10; ++i) { - // expected-error@+3{{expected ']'}} - // expected-note@+2{{to match this '['}} - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+2{{expected ']'}} + // expected-note@+1{{to match this '['}} #pragma acc cache(Arrs.MemArr[3:4:].array[4]) } for (int i = 0; i < 10; ++i) { - // expected-error@+2{{OpenACC sub-array is not allowed here}} - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+1{{OpenACC sub-array is not allowed here}} #pragma acc cache(Arrs.MemArr[:].array[4]) } for (int i = 0; i < 10; ++i) { - // expected-error@+2{{expected unqualified-id}} - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+1{{expected unqualified-id}} #pragma acc cache(Arrs.MemArr[::].array[4]) } for (int i = 0; i < 10; ++i) { - // expected-error@+4{{expected expression}} - // expected-error@+3{{expected ']'}} - // expected-note@+2{{to match this '['}} - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+3{{expected expression}} + // expected-error@+2{{expected ']'}} + // expected-note@+1{{to match this '['}} #pragma acc cache(Arrs.MemArr[: :].array[4]) } for (int i = 0; i < 10; ++i) { - // expected-error@+2{{OpenACC sub-array is not allowed here}} - // expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}} + // expected-error@+1{{OpenACC sub-array is not allowed here}} #pragma acc cache(Arrs.MemArr[3:].array[4]) } - func<S, 5>(); + func<S, 5>(); // expected-note{{in instantiation of function template specialization}} } diff --git a/clang/test/SemaOpenACC/cache-construct-ast.cpp b/clang/test/SemaOpenACC/cache-construct-ast.cpp new file mode 100644 index 0000000000000..3d2e6a0eb0c3f --- /dev/null +++ b/clang/test/SemaOpenACC/cache-construct-ast.cpp @@ -0,0 +1,121 @@ +// RUN: %clang_cc1 %s -fopenacc -ast-dump | FileCheck %s + +// Test this with PCH. +// RUN: %clang_cc1 %s -fopenacc -emit-pch -o %t %s +// RUN: %clang_cc1 %s -fopenacc -include-pch %t -ast-dump-all | FileCheck %s + +#ifndef PCH_HELPER +#define PCH_HELPER + +void use() { + // CHECK: FunctionDecl{{.*}} use 'void ()' + // CHECK-NEXT: CompoundStmt + int Array[5]; + // CHECK-NEXT: DeclStmt + // CHECK-NEXT: VarDecl{{.*}}Array 'int[5]' + +#pragma acc cache(Array[1]) + // CHECK-NEXT: OpenACCCacheConstruct{{.*}} cache + // CHECK-NEXT: ArraySubscriptExpr{{.*}}'int' lvalue + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int *' <ArrayToPointerDecay> + // CHECK-NEXT: DeclRefExpr{{.*}}'Array' 'int[5]' + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1 +#pragma acc cache(Array[1:2]) + // CHECK-NEXT: OpenACCCacheConstruct{{.*}} cache + // CHECK-NEXT: ArraySectionExpr{{.*}}'<array section type>' lvalue + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int *' <ArrayToPointerDecay> + // CHECK-NEXT: DeclRefExpr{{.*}}'Array' 'int[5]' + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1 + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 2 +} + +struct S { + int Array[5]; + int Array2D[5][5]; + + void StructUse() { + // CHECK: CXXMethodDecl{{.*}}StructUse 'void ()' + // CHECK-NEXT: CompoundStmt +#pragma acc cache(Array[1]) + // CHECK-NEXT: OpenACCCacheConstruct{{.*}} cache + // CHECK-NEXT: ArraySubscriptExpr{{.*}}'int' lvalue + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int *' <ArrayToPointerDecay> + // CHECK-NEXT: MemberExpr{{.*}} 'int[5]' lvalue ->Array + // CHECK-NEXT: CXXThisExpr{{.*}} 'S *' + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1 +#pragma acc cache(Array[1:2]) + // CHECK-NEXT: OpenACCCacheConstruct{{.*}} cache + // CHECK-NEXT: ArraySectionExpr{{.*}}'<array section type>' lvalue + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int *' <ArrayToPointerDecay> + // CHECK-NEXT: MemberExpr{{.*}} 'int[5]' lvalue ->Array + // CHECK-NEXT: CXXThisExpr{{.*}} 'S *' + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1 + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 2 +#pragma acc cache(Array2D[1][1]) + // CHECK-NEXT: OpenACCCacheConstruct{{.*}} cache + // CHECK-NEXT: ArraySubscriptExpr{{.*}}'int' lvalue + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int *' <ArrayToPointerDecay> + // CHECK-NEXT: ArraySubscriptExpr{{.*}}'int[5]' lvalue + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int (*)[5]' <ArrayToPointerDecay> + // CHECK-NEXT: MemberExpr{{.*}} 'int[5][5]' lvalue ->Array2D + // CHECK-NEXT: CXXThisExpr{{.*}} 'S *' + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1 + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1 +#pragma acc cache(Array2D[1][1:2]) + // CHECK-NEXT: OpenACCCacheConstruct{{.*}} cache + // CHECK-NEXT: ArraySectionExpr{{.*}}'<array section type>' lvalue + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int *' <ArrayToPointerDecay> + // CHECK-NEXT: ArraySubscriptExpr{{.*}}'int[5]' lvalue + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int (*)[5]' <ArrayToPointerDecay> + // CHECK-NEXT: MemberExpr{{.*}} 'int[5][5]' lvalue ->Array2D + // CHECK-NEXT: CXXThisExpr{{.*}} 'S *' + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1 + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1 + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 2 + } +}; + +template<typename T> +void templ_use() { + // CHECK: FunctionDecl{{.*}} templ_use 'void ()' + // CHECK-NEXT: CompoundStmt + T Array[5]; + // CHECK-NEXT: DeclStmt + // CHECK-NEXT: VarDecl{{.*}}Array 'T[5]' + +#pragma acc cache(Array[1]) + // CHECK-NEXT: OpenACCCacheConstruct{{.*}} cache + // CHECK-NEXT: ArraySubscriptExpr{{.*}}'T' lvalue + // CHECK-NEXT: DeclRefExpr{{.*}}'Array' 'T[5]' + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1 +#pragma acc cache(Array[1:2]) + // CHECK-NEXT: OpenACCCacheConstruct{{.*}} cache + // CHECK-NEXT: ArraySectionExpr{{.*}}'<dependent type>' lvalue + // CHECK-NEXT: DeclRefExpr{{.*}}'Array' 'T[5]' + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1 + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 2 + + // Instantiation: + // CHECK: FunctionDecl{{.*}} templ_use 'void ()' implicit_instantiation + // CHECK-NEXT: TemplateArgument type 'int' + // CHECK-NEXT: BuiltinType{{.*}} 'int' + // CHECK-NEXT: CompoundStmt + // CHECK-NEXT: DeclStmt + // CHECK-NEXT: VarDecl{{.*}}Array 'int[5]' + // CHECK-NEXT: OpenACCCacheConstruct{{.*}} cache + // CHECK-NEXT: ArraySubscriptExpr{{.*}}'int' lvalue + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int *' <ArrayToPointerDecay> + // CHECK-NEXT: DeclRefExpr{{.*}}'Array' 'int[5]' + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1 + // CHECK-NEXT: OpenACCCacheConstruct{{.*}} cache + // CHECK-NEXT: ArraySectionExpr{{.*}}'<array section type>' lvalue + // CHECK-NEXT: ImplicitCastExpr{{.*}}'int *' <ArrayToPointerDecay> + // CHECK-NEXT: DeclRefExpr{{.*}}'Array' 'int[5]' + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1 + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 2 +} + +void foo() { + templ_use<int>(); +} +#endif diff --git a/clang/test/SemaOpenACC/cache-construct.cpp b/clang/test/SemaOpenACC/cache-construct.cpp new file mode 100644 index 0000000000000..0a32754c82710 --- /dev/null +++ b/clang/test/SemaOpenACC/cache-construct.cpp @@ -0,0 +1,51 @@ +// RUN: %clang_cc1 %s -fopenacc -verify + + +void use() { + int Array[5]; + int NotArray; + +#pragma acc cache(Array[1]) +#pragma acc cache(Array[1:2]) + + // expected-error@+1{{OpenACC variable in cache directive is not a valid sub-array or array element}} +#pragma acc cache(Array) + // expected-error@+1{{OpenACC variable in cache directive is not a valid sub-array or array element}} +#pragma acc cache(NotArray) +} + +struct S { + int Array[5]; + int NotArray; + int Array2D[5][5]; + + void use() { +#pragma acc cache(Array[1]) +#pragma acc cache(Array[1:2]) +#pragma acc cache(Array2D[1][1]) +#pragma acc cache(Array2D[1][1:2]) + + // expected-error@+1{{OpenACC variable in cache directive is not a valid sub-array or array element}} +#pragma acc cache(Array) + // expected-error@+1{{OpenACC variable in cache directive is not a valid sub-array or array element}} +#pragma acc cache(NotArray) + } +}; + +template<typename T> +void templ_use() { + T Array[5]; + T NotArray; + +#pragma acc cache(Array[1]) +#pragma acc cache(Array[1:2]) + + // expected-error@+1{{OpenACC variable in cache directive is not a valid sub-array or array element}} +#pragma acc cache(Array) + // expected-error@+1{{OpenACC variable in cache directive is not a valid sub-array or array element}} +#pragma acc cache(NotArray) +} + +void foo() { + templ_use<int>(); +} diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp index 479490442f3c8..dba5516cec88c 100644 --- a/clang/tools/libclang/CIndex.cpp +++ b/clang/tools/libclang/CIndex.cpp @@ -2190,6 +2190,7 @@ class EnqueueVisitor : public ConstStmtVisitor<EnqueueVisitor, void>, void VisitOpenACCExitDataConstruct(const OpenACCExitDataConstruct *D); void VisitOpenACCHostDataConstruct(const OpenACCHostDataConstruct *D); void VisitOpenACCWaitConstruct(const OpenACCWaitConstruct *D); + void VisitOpenACCCacheConstruct(const OpenACCCacheConstruct *D); void VisitOpenACCInitConstruct(const OpenACCInitConstruct *D); void VisitOpenACCShutdownConstruct(const OpenACCShutdownConstruct *D); void VisitOpenACCSetConstruct(const OpenACCSetConstruct *D); @@ -3680,6 +3681,11 @@ void EnqueueVisitor::VisitOpenACCWaitConstruct(const OpenACCWaitConstruct *C) { EnqueueChildren(Clause); } +void EnqueueVisitor::VisitOpenACCCacheConstruct( + const OpenACCCacheConstruct *C) { + EnqueueChildren(C); +} + void EnqueueVisitor::VisitOpenACCInitConstruct(const OpenACCInitConstruct *C) { EnqueueChildren(C); for (auto *Clause : C->clauses()) @@ -6477,6 +6483,8 @@ CXString clang_getCursorKindSpelling(enum CXCursorKind Kind) { return cxstring::createRef("OpenACCHostDataConstruct"); case CXCursor_OpenACCWaitConstruct: return cxstring::createRef("OpenACCWaitConstruct"); + case CXCursor_OpenACCCacheConstruct: + return cxstring::createRef("OpenACCCacheConstruct"); case CXCursor_OpenACCInitConstruct: return cxstring::createRef("OpenACCInitConstruct"); case CXCursor_OpenACCShutdownConstruct: diff --git a/clang/tools/libclang/CXCursor.cpp b/clang/tools/libclang/CXCursor.cpp index 0810c38bb751b..1d15120106017 100644 --- a/clang/tools/libclang/CXCursor.cpp +++ b/clang/tools/libclang/CXCursor.cpp @@ -910,6 +910,9 @@ CXCursor cxcursor::MakeCXCursor(const Stmt *S, const Decl *Parent, case Stmt::OpenACCWaitConstructClass: K = CXCursor_OpenACCWaitConstruct; break; + case Stmt::OpenACCCacheConstructClass: + K = CXCursor_OpenACCCacheConstruct; + break; case Stmt::OpenACCInitConstructClass: K = CXCursor_OpenACCInitConstruct; break; _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
