[clang] d5cec38 - [OpenACC] Implement 'cache' construct AST/Sema
via cfe-commits
cfe-commits at lists.llvm.org
Mon Mar 3 13:57:27 PST 2025
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 at +2{{expected '('}}
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +1{{expected '('}}
#pragma acc cache
}
for (int i = 0; i < 10; ++i) {
- // expected-error at +3{{expected '('}}
- // expected-error at +2{{invalid OpenACC clause 'clause'}}
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +2{{expected '('}}
+ // expected-error at +1{{invalid OpenACC clause 'clause'}}
#pragma acc cache clause list
}
for (int i = 0; i < 10; ++i) {
- // expected-error at +2{{expected expression}}
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +1{{expected expression}}
#pragma acc cache()
}
for (int i = 0; i < 10; ++i) {
- // expected-error at +3{{expected expression}}
- // expected-error at +2{{invalid OpenACC clause 'clause'}}
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +2{{expected expression}}
+ // expected-error at +1{{invalid OpenACC clause 'clause'}}
#pragma acc cache() clause-list
}
for (int i = 0; i < 10; ++i) {
- // expected-error at +3{{expected ')'}}
- // expected-note at +2{{to match this '('}}
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +2{{expected ')'}}
+ // expected-note at +1{{to match this '('}}
#pragma acc cache(
}
for (int i = 0; i < 10; ++i) {
- // expected-error at +4{{use of undeclared identifier 'invalid'}}
- // expected-error at +3{{expected ')'}}
- // expected-note at +2{{to match this '('}}
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +3{{use of undeclared identifier 'invalid'}}
+ // expected-error at +2{{expected ')'}}
+ // expected-note at +1{{to match this '('}}
#pragma acc cache(invalid
}
for (int i = 0; i < 10; ++i) {
// expected-error at +3{{expected ')'}}
// expected-note at +2{{to match this '('}}
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +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 at +2{{use of undeclared identifier 'invalid'}}
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +1{{use of undeclared identifier 'invalid'}}
#pragma acc cache(invalid)
}
for (int i = 0; i < 10; ++i) {
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +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 at +6{{expected expression}}
- // expected-error at +5{{expected ']'}}
- // expected-note at +4{{to match this '['}}
- // expected-error at +3{{expected ')'}}
- // expected-note at +2{{to match this '('}}
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +5{{expected expression}}
+ // expected-error at +4{{expected ']'}}
+ // expected-note at +3{{to match this '['}}
+ // expected-error at +2{{expected ')'}}
+ // expected-note at +1{{to match this '('}}
#pragma acc cache(ArrayPtr[
}
for (int i = 0; i < 10; ++i) {
- // expected-error at +4{{expected expression}}
- // expected-error at +3{{expected ']'}}
- // expected-note at +2{{to match this '['}}
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +3{{expected expression}}
+ // expected-error at +2{{expected ']'}}
+ // expected-note at +1{{to match this '['}}
#pragma acc cache(ArrayPtr[, 5)
}
for (int i = 0; i < 10; ++i) {
- // expected-error at +4{{expected expression}}
- // expected-error at +3{{expected ']'}}
- // expected-note at +2{{to match this '['}}
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +3{{expected expression}}
+ // expected-error at +2{{expected ']'}}
+ // expected-note at +1{{to match this '['}}
#pragma acc cache(Array[)
}
for (int i = 0; i < 10; ++i) {
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
#pragma acc cache(Array[*readonly])
}
for (int i = 0; i < 10; ++i) {
- // expected-error at +6{{expected expression}}
- // expected-error at +5{{expected ']'}}
- // expected-note at +4{{to match this '['}}
- // expected-error at +3{{expected ')'}}
- // expected-note at +2{{to match this '('}}
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +5{{expected expression}}
+ // expected-error at +4{{expected ']'}}
+ // expected-note at +3{{to match this '['}}
+ // expected-error at +2{{expected ')'}}
+ // expected-note at +1{{to match this '('}}
#pragma acc cache(Array[*readonly:
}
for (int i = 0; i < 10; ++i) {
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +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 at +2{{invalid tag 'devnum' on 'cache' directive}}
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +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 at +2{{invalid tag 'invalid' on 'cache' directive}}
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +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 at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +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 at +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 at +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 at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +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 at +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 at +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 at +4{{expected expression}}
- // expected-error at +3{{expected ')'}}
- // expected-note at +2{{to match this '('}}
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +3{{expected expression}}
+ // expected-error at +2{{expected ')'}}
+ // expected-note at +1{{to match this '('}}
#pragma acc cache(readonly:ArrayPtr[5:*readonly],
}
for (int i = 0; i < 10; ++i) {
- // expected-error at +2{{expected expression}}
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +1{{expected expression}}
#pragma acc cache(readonly:ArrayPtr[5:*readonly],)
}
for (int i = 0; i < 10; ++i) {
- // expected-warning at +2{{left operand of comma operator has no effect}}
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-warning at +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 at +2{{left operand of comma operator has no effect}}
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-warning at +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 at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +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 at +2{{left operand of comma operator has no effect}}
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-warning at +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 at +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 at +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 at +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 at +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 at +2{{use of undeclared identifier 'NSArray'; did you mean 'NS::NSArray'}}
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +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 at +2{{use of undeclared identifier 'NSInt'; did you mean 'NS::NSInt'}}
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +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 at +2{{use of undeclared identifier 'NSInt'; did you mean 'NS::NSInt'}}
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +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 at +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 at +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 at +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 at +2{{OpenACC sub-array is not allowed here}}
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +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 at +2{{OpenACC sub-array is not allowed here}}
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +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 at +3{{expected ']'}}
- // expected-note at +2{{to match this '['}}
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +2{{expected ']'}}
+ // expected-note at +1{{to match this '['}}
#pragma acc cache(Arrs.MemArr[3:4:].array[4])
}
for (int i = 0; i < 10; ++i) {
- // expected-error at +2{{OpenACC sub-array is not allowed here}}
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +1{{OpenACC sub-array is not allowed here}}
#pragma acc cache(Arrs.MemArr[:].array[4])
}
for (int i = 0; i < 10; ++i) {
- // expected-error at +2{{expected unqualified-id}}
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +1{{expected unqualified-id}}
#pragma acc cache(Arrs.MemArr[::].array[4])
}
for (int i = 0; i < 10; ++i) {
- // expected-error at +4{{expected expression}}
- // expected-error at +3{{expected ']'}}
- // expected-note at +2{{to match this '['}}
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +3{{expected expression}}
+ // expected-error at +2{{expected ']'}}
+ // expected-note at +1{{to match this '['}}
#pragma acc cache(Arrs.MemArr[: :].array[4])
}
for (int i = 0; i < 10; ++i) {
- // expected-error at +2{{OpenACC sub-array is not allowed here}}
- // expected-warning at +1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
+ // expected-error at +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 at +1{{OpenACC variable in cache directive is not a valid sub-array or array element}}
+#pragma acc cache(Array)
+ // expected-error at +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 at +1{{OpenACC variable in cache directive is not a valid sub-array or array element}}
+#pragma acc cache(Array)
+ // expected-error at +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 at +1{{OpenACC variable in cache directive is not a valid sub-array or array element}}
+#pragma acc cache(Array)
+ // expected-error at +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;
More information about the cfe-commits
mailing list