[clang] 5d7d66b - [OpenACC] Implement 'declare' construct AST/Sema
via cfe-commits
cfe-commits at lists.llvm.org
Mon Mar 3 07:48:35 PST 2025
Author: erichkeane
Date: 2025-03-03T07:48:29-08:00
New Revision: 5d7d66ba0d1ad6fcf1aefffd045eea88597f4614
URL: https://github.com/llvm/llvm-project/commit/5d7d66ba0d1ad6fcf1aefffd045eea88597f4614
DIFF: https://github.com/llvm/llvm-project/commit/5d7d66ba0d1ad6fcf1aefffd045eea88597f4614.diff
LOG: [OpenACC] Implement 'declare' construct AST/Sema
The 'declare' construct is the first of two 'declaration' level
constructs, so it is legal in any place a declaration is, including as a
statement, which this accomplishes by wrapping it in a DeclStmt. All
clauses on this have a 'same scope' requirement, which this enforces as
declaration context instead, which makes it possible to implement these
as a template.
The 'link' and 'device_resident' clauses are also added, which have some
similar/small restrictions, but are otherwise pretty rote.
This patch implements all of the above.
Added:
clang/include/clang/AST/DeclOpenACC.h
clang/lib/AST/DeclOpenACC.cpp
clang/test/AST/ast-print-openacc-declare-construct.cpp
clang/test/SemaOpenACC/declare-construct-ast.cpp
clang/test/SemaOpenACC/declare-construct.cpp
Modified:
clang/include/clang/AST/DeclVisitor.h
clang/include/clang/AST/JSONNodeDumper.h
clang/include/clang/AST/OpenACCClause.h
clang/include/clang/AST/RecursiveASTVisitor.h
clang/include/clang/AST/TextNodeDumper.h
clang/include/clang/Basic/DeclNodes.td
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/include/clang/Basic/OpenACCClauses.def
clang/include/clang/Parse/Parser.h
clang/include/clang/Sema/SemaBase.h
clang/include/clang/Sema/SemaOpenACC.h
clang/include/clang/Serialization/ASTBitCodes.h
clang/include/clang/Serialization/ASTRecordReader.h
clang/lib/AST/ASTStructuralEquivalence.cpp
clang/lib/AST/CMakeLists.txt
clang/lib/AST/DeclBase.cpp
clang/lib/AST/DeclPrinter.cpp
clang/lib/AST/JSONNodeDumper.cpp
clang/lib/AST/OpenACCClause.cpp
clang/lib/AST/StmtPrinter.cpp
clang/lib/AST/StmtProfile.cpp
clang/lib/AST/TextNodeDumper.cpp
clang/lib/CodeGen/CGDecl.cpp
clang/lib/CodeGen/CodeGenModule.h
clang/lib/Parse/ParseOpenACC.cpp
clang/lib/Sema/SemaBase.cpp
clang/lib/Sema/SemaOpenACC.cpp
clang/lib/Sema/SemaOpenACCClause.cpp
clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
clang/lib/Sema/TreeTransform.h
clang/lib/Serialization/ASTCommon.cpp
clang/lib/Serialization/ASTReader.cpp
clang/lib/Serialization/ASTReaderDecl.cpp
clang/lib/Serialization/ASTWriter.cpp
clang/lib/Serialization/ASTWriterDecl.cpp
clang/test/ParserOpenACC/parse-clauses.c
clang/test/ParserOpenACC/parse-constructs.c
clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
clang/test/SemaOpenACC/combined-construct-device_type-clause.c
clang/test/SemaOpenACC/compute-construct-device_type-clause.c
clang/test/SemaOpenACC/data-construct.cpp
clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
clang/test/SemaOpenACC/loop-construct-device_type-clause.c
clang/test/SemaOpenACC/unimplemented-construct.c
clang/tools/libclang/CIndex.cpp
Removed:
################################################################################
diff --git a/clang/include/clang/AST/DeclOpenACC.h b/clang/include/clang/AST/DeclOpenACC.h
new file mode 100644
index 0000000000000..fdeb1315cfd09
--- /dev/null
+++ b/clang/include/clang/AST/DeclOpenACC.h
@@ -0,0 +1,106 @@
+//=- DeclOpenACC.h - Classes for representing OpenACC directives -*- C++ -*-==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// This file defines OpenACC nodes for declarative directives.
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_CLANG_AST_DECLOPENACC_H
+#define LLVM_CLANG_AST_DECLOPENACC_H
+
+#include "clang/AST/ASTContext.h"
+#include "clang/AST/Decl.h"
+#include "clang/AST/OpenACCClause.h"
+#include "clang/Basic/OpenACCKinds.h"
+
+namespace clang {
+
+// A base class for the declaration constructs, which manages the clauses and
+// basic source location information. Currently not part of the Decl inheritence
+// tree, as we should never have a reason to store one of these.
+class OpenACCConstructDecl : public Decl {
+ friend class ASTDeclReader;
+ friend class ASTDeclWriter;
+ // The directive kind, each implementation of this interface is expected to
+ // handle a specific kind.
+ OpenACCDirectiveKind DirKind = OpenACCDirectiveKind::Invalid;
+ SourceLocation DirectiveLoc;
+ SourceLocation EndLoc;
+ /// The list of clauses. This is stored here as an ArrayRef, as this is the
+ /// most convienient place to access the list, however the list itself should
+ /// be stored in leaf nodes, likely in trailing-storage.
+ MutableArrayRef<const OpenACCClause *> Clauses;
+
+protected:
+ OpenACCConstructDecl(Kind DeclKind, DeclContext *DC, OpenACCDirectiveKind K,
+ SourceLocation StartLoc, SourceLocation DirLoc,
+ SourceLocation EndLoc)
+ : Decl(DeclKind, DC, StartLoc), DirKind(K), DirectiveLoc(DirLoc),
+ EndLoc(EndLoc) {}
+
+ OpenACCConstructDecl(Kind DeclKind) : Decl(DeclKind, EmptyShell{}) {}
+
+ void setClauseList(MutableArrayRef<const OpenACCClause *> NewClauses) {
+ assert(Clauses.empty() && "Cannot change clause list");
+ Clauses = NewClauses;
+ }
+
+public:
+ OpenACCDirectiveKind getDirectiveKind() const { return DirKind; }
+ SourceLocation getDirectiveLoc() const { return DirectiveLoc; }
+ virtual SourceRange getSourceRange() const override LLVM_READONLY {
+ return SourceRange(getLocation(), EndLoc);
+ }
+
+ ArrayRef<const OpenACCClause *> clauses() const { return Clauses; }
+};
+
+class OpenACCDeclareDecl final
+ : public OpenACCConstructDecl,
+ private llvm::TrailingObjects<OpenACCDeclareDecl, const OpenACCClause *> {
+ friend TrailingObjects;
+ friend class ASTDeclReader;
+ friend class ASTDeclWriter;
+
+ OpenACCDeclareDecl(unsigned NumClauses)
+ : OpenACCConstructDecl(OpenACCDeclare) {
+ std::uninitialized_value_construct(
+ getTrailingObjects<const OpenACCClause *>(),
+ getTrailingObjects<const OpenACCClause *>() + NumClauses);
+ setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
+ NumClauses));
+ }
+
+ OpenACCDeclareDecl(DeclContext *DC, SourceLocation StartLoc,
+ SourceLocation DirLoc, SourceLocation EndLoc,
+ ArrayRef<const OpenACCClause *> Clauses)
+ : OpenACCConstructDecl(OpenACCDeclare, DC, OpenACCDirectiveKind::Declare,
+ StartLoc, DirLoc, EndLoc) {
+ // Initialize the trailing storage.
+ std::uninitialized_copy(Clauses.begin(), Clauses.end(),
+ getTrailingObjects<const OpenACCClause *>());
+
+ setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
+ Clauses.size()));
+ }
+
+public:
+ static OpenACCDeclareDecl *Create(ASTContext &Ctx, DeclContext *DC,
+ SourceLocation StartLoc,
+ SourceLocation DirLoc,
+ SourceLocation EndLoc,
+ ArrayRef<const OpenACCClause *> Clauses);
+ static OpenACCDeclareDecl *
+ CreateDeserialized(ASTContext &Ctx, GlobalDeclID ID, unsigned NumClauses);
+ static bool classof(const Decl *D) { return classofKind(D->getKind()); }
+ static bool classofKind(Kind K) { return K == OpenACCDeclare; }
+};
+} // namespace clang
+
+#endif
diff --git a/clang/include/clang/AST/DeclVisitor.h b/clang/include/clang/AST/DeclVisitor.h
index 8690cdda4bb70..8b7c30ee4051e 100644
--- a/clang/include/clang/AST/DeclVisitor.h
+++ b/clang/include/clang/AST/DeclVisitor.h
@@ -18,6 +18,7 @@
#include "clang/AST/DeclCXX.h"
#include "clang/AST/DeclFriend.h"
#include "clang/AST/DeclObjC.h"
+#include "clang/AST/DeclOpenACC.h"
#include "clang/AST/DeclOpenMP.h"
#include "clang/AST/DeclTemplate.h"
#include "llvm/ADT/STLExtras.h"
diff --git a/clang/include/clang/AST/JSONNodeDumper.h b/clang/include/clang/AST/JSONNodeDumper.h
index 660a05435003d..97b07e1afbf68 100644
--- a/clang/include/clang/AST/JSONNodeDumper.h
+++ b/clang/include/clang/AST/JSONNodeDumper.h
@@ -281,6 +281,8 @@ class JSONNodeDumper
void VisitObjCPropertyImplDecl(const ObjCPropertyImplDecl *D);
void VisitBlockDecl(const BlockDecl *D);
+ void VisitOpenACCDeclareDecl(const OpenACCDeclareDecl *D);
+
void VisitDeclRefExpr(const DeclRefExpr *DRE);
void VisitSYCLUniqueStableNameExpr(const SYCLUniqueStableNameExpr *E);
void VisitOpenACCAsteriskSizeExpr(const OpenACCAsteriskSizeExpr *E);
diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index f5be54bdada8b..0fc638b73bbd2 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -1171,6 +1171,55 @@ class OpenACCReductionClause final
OpenACCReductionOperator getReductionOp() const { return Op; }
};
+class OpenACCLinkClause final
+ : public OpenACCClauseWithVarList,
+ private llvm::TrailingObjects<OpenACCLinkClause, Expr *> {
+ friend TrailingObjects;
+
+ OpenACCLinkClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
+ ArrayRef<Expr *> VarList, SourceLocation EndLoc)
+ : OpenACCClauseWithVarList(OpenACCClauseKind::Link, BeginLoc, LParenLoc,
+ EndLoc) {
+ std::uninitialized_copy(VarList.begin(), VarList.end(),
+ getTrailingObjects<Expr *>());
+ setExprs(MutableArrayRef(getTrailingObjects<Expr *>(), VarList.size()));
+ }
+
+public:
+ static bool classof(const OpenACCClause *C) {
+ return C->getClauseKind() == OpenACCClauseKind::Link;
+ }
+
+ static OpenACCLinkClause *Create(const ASTContext &C, SourceLocation BeginLoc,
+ SourceLocation LParenLoc,
+ ArrayRef<Expr *> VarList,
+ SourceLocation EndLoc);
+};
+
+class OpenACCDeviceResidentClause final
+ : public OpenACCClauseWithVarList,
+ private llvm::TrailingObjects<OpenACCDeviceResidentClause, Expr *> {
+ friend TrailingObjects;
+
+ OpenACCDeviceResidentClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
+ ArrayRef<Expr *> VarList, SourceLocation EndLoc)
+ : OpenACCClauseWithVarList(OpenACCClauseKind::DeviceResident, BeginLoc,
+ LParenLoc, EndLoc) {
+ std::uninitialized_copy(VarList.begin(), VarList.end(),
+ getTrailingObjects<Expr *>());
+ setExprs(MutableArrayRef(getTrailingObjects<Expr *>(), VarList.size()));
+ }
+
+public:
+ static bool classof(const OpenACCClause *C) {
+ return C->getClauseKind() == OpenACCClauseKind::DeviceResident;
+ }
+
+ static OpenACCDeviceResidentClause *
+ Create(const ASTContext &C, SourceLocation BeginLoc, SourceLocation LParenLoc,
+ ArrayRef<Expr *> VarList, SourceLocation EndLoc);
+};
+
template <class Impl> class OpenACCClauseVisitor {
Impl &getDerived() { return static_cast<Impl &>(*this); }
diff --git a/clang/include/clang/AST/RecursiveASTVisitor.h b/clang/include/clang/AST/RecursiveASTVisitor.h
index 5964cbaec8e44..b1b4363b65721 100644
--- a/clang/include/clang/AST/RecursiveASTVisitor.h
+++ b/clang/include/clang/AST/RecursiveASTVisitor.h
@@ -20,6 +20,7 @@
#include "clang/AST/DeclCXX.h"
#include "clang/AST/DeclFriend.h"
#include "clang/AST/DeclObjC.h"
+#include "clang/AST/DeclOpenACC.h"
#include "clang/AST/DeclOpenMP.h"
#include "clang/AST/DeclTemplate.h"
#include "clang/AST/DeclarationName.h"
@@ -1821,6 +1822,9 @@ DEF_TRAVERSE_DECL(OMPAllocateDecl, {
TRY_TO(TraverseOMPClause(C));
})
+DEF_TRAVERSE_DECL(OpenACCDeclareDecl,
+ { TRY_TO(VisitOpenACCClauseList(D->clauses())); })
+
// A helper method for TemplateDecl's children.
template <typename Derived>
bool RecursiveASTVisitor<Derived>::TraverseTemplateParameterListHelper(
diff --git a/clang/include/clang/AST/TextNodeDumper.h b/clang/include/clang/AST/TextNodeDumper.h
index 81844db2c77fa..b0e5b28e22863 100644
--- a/clang/include/clang/AST/TextNodeDumper.h
+++ b/clang/include/clang/AST/TextNodeDumper.h
@@ -423,6 +423,7 @@ class TextNodeDumper
void VisitOpenACCUpdateConstruct(const OpenACCUpdateConstruct *S);
void VisitOpenACCAtomicConstruct(const OpenACCAtomicConstruct *S);
void VisitOpenACCAsteriskSizeExpr(const OpenACCAsteriskSizeExpr *S);
+ void VisitOpenACCDeclareDecl(const OpenACCDeclareDecl *D);
void VisitEmbedExpr(const EmbedExpr *S);
void VisitAtomicExpr(const AtomicExpr *AE);
void VisitConvertVectorExpr(const ConvertVectorExpr *S);
diff --git a/clang/include/clang/Basic/DeclNodes.td b/clang/include/clang/Basic/DeclNodes.td
index 723113dc2486e..227184dc71cc7 100644
--- a/clang/include/clang/Basic/DeclNodes.td
+++ b/clang/include/clang/Basic/DeclNodes.td
@@ -111,3 +111,4 @@ def Empty : DeclNode<Decl>;
def RequiresExprBody : DeclNode<Decl>, DeclContext;
def LifetimeExtendedTemporary : DeclNode<Decl>;
def HLSLBuffer : DeclNode<Named, "HLSLBuffer">, DeclContext;
+def OpenACCDeclare : DeclNode<Decl, "#pragma acc declare">;
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index d094c075ecee2..0efb15405ed5d 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12812,6 +12812,7 @@ def err_acc_duplicate_clause_disallowed
"directive">;
def note_acc_previous_clause_here : Note<"previous clause is here">;
def note_acc_previous_expr_here : Note<"previous expression is here">;
+def note_acc_previous_reference : Note<"previous reference is here">;
def err_acc_branch_in_out_compute_construct
: Error<"invalid %select{branch|return|throw}0 %select{out of|into}1 "
"OpenACC Compute/Combined Construct">;
@@ -12844,9 +12845,9 @@ def err_acc_not_a_var_ref
: Error<"OpenACC variable is not a valid variable name, sub-array, array "
"element,%select{| member of a composite variable,}0 or composite "
"variable member">;
-def err_acc_not_a_var_ref_use_device
- : Error<"OpenACC variable in 'use_device' clause is not a valid variable "
- "name or array name">;
+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_typecheck_subarray_value
: Error<"OpenACC sub-array subscripted value is not an array or pointer">;
def err_acc_subarray_function_type
@@ -13014,6 +13015,23 @@ def note_acc_atomic_mismatch_compound_operand
"side of assignment|<not possible>|on left hand side of compound "
"assignment|on left hand side of assignment}0('%1') from the first "
"statement">;
+def err_acc_declare_required_clauses
+ : Error<"no valid clauses specified in OpenACC 'declare' directive">;
+def err_acc_declare_clause_at_global
+ : Error<"OpenACC '%0' clause on a 'declare' directive is not allowed at "
+ "global or namespace scope">;
+def err_acc_link_not_extern
+ : Error<"variable referenced by 'link' clause not in global or namespace "
+ "scope must be marked 'extern'">;
+def err_acc_declare_extern
+ : Error<"'extern' variable may not be referenced by '%0' clause on an "
+ "OpenACC 'declare' directive">;
+def err_acc_declare_same_scope
+ : Error<"variable appearing in '%0' clause of OpenACC 'declare' directive "
+ "must be in the same scope as the directive">;
+def err_acc_multiple_references
+ : Error<"variable referenced in '%0' clause of OpenACC 'declare' directive "
+ "was already referenced">;
// AMDGCN builtins diagnostics
def err_amdgcn_global_load_lds_size_invalid_value : Error<"invalid size value">;
diff --git a/clang/include/clang/Basic/OpenACCClauses.def b/clang/include/clang/Basic/OpenACCClauses.def
index 8b15007c85557..1d9e3b09a749b 100644
--- a/clang/include/clang/Basic/OpenACCClauses.def
+++ b/clang/include/clang/Basic/OpenACCClauses.def
@@ -44,6 +44,7 @@ VISIT_CLAUSE(Detach)
VISIT_CLAUSE(Device)
VISIT_CLAUSE(DeviceNum)
VISIT_CLAUSE(DevicePtr)
+VISIT_CLAUSE(DeviceResident)
VISIT_CLAUSE(DeviceType)
CLAUSE_ALIAS(DType, DeviceType, false)
VISIT_CLAUSE(Finalize)
@@ -53,6 +54,7 @@ VISIT_CLAUSE(Host)
VISIT_CLAUSE(If)
VISIT_CLAUSE(IfPresent)
VISIT_CLAUSE(Independent)
+VISIT_CLAUSE(Link)
VISIT_CLAUSE(NoCreate)
VISIT_CLAUSE(NumGangs)
VISIT_CLAUSE(NumWorkers)
diff --git a/clang/include/clang/Parse/Parser.h b/clang/include/clang/Parse/Parser.h
index 335258d597028..0602f44333b20 100644
--- a/clang/include/clang/Parse/Parser.h
+++ b/clang/include/clang/Parse/Parser.h
@@ -3756,9 +3756,11 @@ class Parser : public CodeCompletionHandler {
using OpenACCVarParseResult = std::pair<ExprResult, OpenACCParseCanContinue>;
/// Parses a single variable in a variable list for OpenACC.
- OpenACCVarParseResult ParseOpenACCVar(OpenACCClauseKind CK);
+ OpenACCVarParseResult ParseOpenACCVar(OpenACCDirectiveKind DK,
+ OpenACCClauseKind CK);
/// Parses the variable list for the variety of places that take a var-list.
- llvm::SmallVector<Expr *> ParseOpenACCVarList(OpenACCClauseKind CK);
+ llvm::SmallVector<Expr *> ParseOpenACCVarList(OpenACCDirectiveKind DK,
+ OpenACCClauseKind CK);
/// Parses any parameters for an OpenACC Clause, including required/optional
/// parens.
OpenACCClauseParseResult
diff --git a/clang/include/clang/Sema/SemaBase.h b/clang/include/clang/Sema/SemaBase.h
index 0b05245ab9686..463cae83c7e81 100644
--- a/clang/include/clang/Sema/SemaBase.h
+++ b/clang/include/clang/Sema/SemaBase.h
@@ -42,6 +42,7 @@ class SemaBase {
ASTContext &getASTContext() const;
DiagnosticsEngine &getDiagnostics() const;
const LangOptions &getLangOpts() const;
+ DeclContext *getCurContext() const;
/// Helper class that creates diagnostics with optional
/// template instantiation stacks.
diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index 3004b98760a98..36a2b6ff7edc3 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -33,6 +33,9 @@ class IdentifierInfo;
class OpenACCClause;
class SemaOpenACC : public SemaBase {
+public:
+ using DeclGroupPtrTy = OpaquePtr<DeclGroupRef>;
+
private:
struct ComputeConstructInfo {
/// Which type of compute construct we are inside of, which we can use to
@@ -158,6 +161,13 @@ class SemaOpenACC : public SemaBase {
/// Helper function for checking the 'for' and 'range for' stmts.
void ForStmtBeginHelper(SourceLocation ForLoc, ForStmtBeginChecker &C);
+ // The 'declare' construct requires only a single reference among ALL declare
+ // directives in a context. We store existing references to check. Because the
+ // rules prevent referencing the same variable from multiple declaration
+ // contexts, we can just store the declaration and location of the reference.
+ llvm::DenseMap<const clang::DeclaratorDecl *, SourceLocation>
+ DeclareVarReferences;
+
public:
ComputeConstructInfo &getActiveComputeConstructInfo() {
return ActiveComputeConstructInfo;
@@ -411,6 +421,8 @@ class SemaOpenACC : public SemaBase {
ClauseKind == OpenACCClauseKind::Reduction ||
ClauseKind == OpenACCClauseKind::Host ||
ClauseKind == OpenACCClauseKind::Device ||
+ ClauseKind == OpenACCClauseKind::DeviceResident ||
+ ClauseKind == OpenACCClauseKind::Link ||
(ClauseKind == OpenACCClauseKind::Self &&
DirKind == OpenACCDirectiveKind::Update) ||
ClauseKind == OpenACCClauseKind::FirstPrivate) &&
@@ -427,23 +439,10 @@ class SemaOpenACC : public SemaBase {
}
bool isReadOnly() const {
- assert((ClauseKind == OpenACCClauseKind::CopyIn ||
- ClauseKind == OpenACCClauseKind::PCopyIn ||
- ClauseKind == OpenACCClauseKind::PresentOrCopyIn) &&
- "Only copyin accepts 'readonly:' tag");
return std::get<VarListDetails>(Details).IsReadOnly;
}
- bool isZero() const {
- assert((ClauseKind == OpenACCClauseKind::CopyOut ||
- ClauseKind == OpenACCClauseKind::PCopyOut ||
- ClauseKind == OpenACCClauseKind::PresentOrCopyOut ||
- ClauseKind == OpenACCClauseKind::Create ||
- ClauseKind == OpenACCClauseKind::PCreate ||
- ClauseKind == OpenACCClauseKind::PresentOrCreate) &&
- "Only copyout/create accepts 'zero' tag");
- return std::get<VarListDetails>(Details).IsZero;
- }
+ bool isZero() const { return std::get<VarListDetails>(Details).IsZero; }
bool isForce() const {
assert(ClauseKind == OpenACCClauseKind::Collapse &&
@@ -557,6 +556,8 @@ class SemaOpenACC : public SemaBase {
ClauseKind == OpenACCClauseKind::DevicePtr ||
ClauseKind == OpenACCClauseKind::Host ||
ClauseKind == OpenACCClauseKind::Device ||
+ ClauseKind == OpenACCClauseKind::DeviceResident ||
+ ClauseKind == OpenACCClauseKind::Link ||
(ClauseKind == OpenACCClauseKind::Self &&
DirKind == OpenACCDirectiveKind::Update) ||
ClauseKind == OpenACCClauseKind::FirstPrivate) &&
@@ -600,6 +601,8 @@ class SemaOpenACC : public SemaBase {
ClauseKind == OpenACCClauseKind::DevicePtr ||
ClauseKind == OpenACCClauseKind::Host ||
ClauseKind == OpenACCClauseKind::Device ||
+ ClauseKind == OpenACCClauseKind::DeviceResident ||
+ ClauseKind == OpenACCClauseKind::Link ||
(ClauseKind == OpenACCClauseKind::Self &&
DirKind == OpenACCDirectiveKind::Update) ||
ClauseKind == OpenACCClauseKind::FirstPrivate) &&
@@ -744,7 +747,11 @@ class SemaOpenACC : public SemaBase {
/// Called after the directive has been completely parsed, including the
/// declaration group or associated statement.
- DeclGroupRef ActOnEndDeclDirective();
+ DeclGroupRef ActOnEndDeclDirective(OpenACCDirectiveKind K,
+ SourceLocation StartLoc,
+ SourceLocation DirLoc,
+ SourceLocation EndLoc,
+ ArrayRef<OpenACCClause *> Clauses);
/// Called when encountering an 'int-expr' for OpenACC, and manages
/// conversions and diagnostics to 'int'.
@@ -753,7 +760,16 @@ class SemaOpenACC : public SemaBase {
/// Called when encountering a 'var' for OpenACC, ensures it is actually a
/// declaration reference to a variable of the correct type.
- ExprResult ActOnVar(OpenACCClauseKind CK, Expr *VarExpr);
+ ExprResult ActOnVar(OpenACCDirectiveKind DK, OpenACCClauseKind CK,
+ Expr *VarExpr);
+
+ // Called after 'ActOnVar' specifically for a 'link' clause, which has to do
+ // some minor additional checks.
+ llvm::SmallVector<Expr *> CheckLinkClauseVarList(ArrayRef<Expr *> VarExpr);
+
+ // Checking for the arguments specific to the declare-clause that need to be
+ // checked during both phases of template translation.
+ bool CheckDeclareClause(SemaOpenACC::OpenACCParsedClause &Clause);
/// Called while semantically analyzing the reduction clause, ensuring the var
/// is the correct kind of reference.
diff --git a/clang/include/clang/Serialization/ASTBitCodes.h b/clang/include/clang/Serialization/ASTBitCodes.h
index c998be34b9d89..5698bebf13445 100644
--- a/clang/include/clang/Serialization/ASTBitCodes.h
+++ b/clang/include/clang/Serialization/ASTBitCodes.h
@@ -1528,7 +1528,10 @@ enum DeclCode {
// A decls specialization record.
DECL_PARTIAL_SPECIALIZATIONS,
- DECL_LAST = DECL_PARTIAL_SPECIALIZATIONS
+ // An OpenACCDeclareDecl record.
+ DECL_OPENACC_DECLARE,
+
+ DECL_LAST = DECL_OPENACC_DECLARE
};
/// Record codes for each kind of statement or expression.
diff --git a/clang/include/clang/Serialization/ASTRecordReader.h b/clang/include/clang/Serialization/ASTRecordReader.h
index 2561418b78ca7..ae7cd84fbc647 100644
--- a/clang/include/clang/Serialization/ASTRecordReader.h
+++ b/clang/include/clang/Serialization/ASTRecordReader.h
@@ -278,7 +278,8 @@ class ASTRecordReader
/// Read an OpenACC clause, advancing Idx.
OpenACCClause *readOpenACCClause();
- /// Read a list of OpenACC clauses into the passed SmallVector.
+ /// Read a list of OpenACC clauses into the passed SmallVector, during
+ /// statement reading.
void readOpenACCClauseList(MutableArrayRef<const OpenACCClause *> Clauses);
/// Read a source location, advancing Idx.
diff --git a/clang/lib/AST/ASTStructuralEquivalence.cpp b/clang/lib/AST/ASTStructuralEquivalence.cpp
index eaf0748395268..b74f67f0a9fed 100644
--- a/clang/lib/AST/ASTStructuralEquivalence.cpp
+++ b/clang/lib/AST/ASTStructuralEquivalence.cpp
@@ -66,6 +66,7 @@
#include "clang/AST/DeclCXX.h"
#include "clang/AST/DeclFriend.h"
#include "clang/AST/DeclObjC.h"
+#include "clang/AST/DeclOpenACC.h"
#include "clang/AST/DeclOpenMP.h"
#include "clang/AST/DeclTemplate.h"
#include "clang/AST/ExprCXX.h"
diff --git a/clang/lib/AST/CMakeLists.txt b/clang/lib/AST/CMakeLists.txt
index cb13c5225b713..f6056e3935a63 100644
--- a/clang/lib/AST/CMakeLists.txt
+++ b/clang/lib/AST/CMakeLists.txt
@@ -50,6 +50,7 @@ add_clang_library(clangAST
DeclFriend.cpp
DeclGroup.cpp
DeclObjC.cpp
+ DeclOpenACC.cpp
DeclOpenMP.cpp
DeclPrinter.cpp
DeclTemplate.cpp
diff --git a/clang/lib/AST/DeclBase.cpp b/clang/lib/AST/DeclBase.cpp
index ab9d4869a74ee..9d8eb07f255ed 100644
--- a/clang/lib/AST/DeclBase.cpp
+++ b/clang/lib/AST/DeclBase.cpp
@@ -21,6 +21,7 @@
#include "clang/AST/DeclContextInternals.h"
#include "clang/AST/DeclFriend.h"
#include "clang/AST/DeclObjC.h"
+#include "clang/AST/DeclOpenACC.h"
#include "clang/AST/DeclOpenMP.h"
#include "clang/AST/DeclTemplate.h"
#include "clang/AST/DependentDiagnostic.h"
@@ -992,6 +993,7 @@ unsigned Decl::getIdentifierNamespaceForKind(Kind DeclKind) {
case LifetimeExtendedTemporary:
case RequiresExprBody:
case ImplicitConceptSpecialization:
+ case OpenACCDeclare:
// Never looked up by name.
return 0;
}
diff --git a/clang/lib/AST/DeclOpenACC.cpp b/clang/lib/AST/DeclOpenACC.cpp
new file mode 100644
index 0000000000000..568fd9dcceec7
--- /dev/null
+++ b/clang/lib/AST/DeclOpenACC.cpp
@@ -0,0 +1,33 @@
+//===--- DeclOpenACC.cpp - Classes for OpenACC Constructs -----------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file implements the subclasses of Decl class declared in Decl.h
+//
+//===----------------------------------------------------------------------===//
+
+#include "clang/AST/DeclOpenACC.h"
+#include "clang/AST/ASTContext.h"
+
+using namespace clang;
+
+OpenACCDeclareDecl *
+OpenACCDeclareDecl::Create(ASTContext &Ctx, DeclContext *DC,
+ SourceLocation StartLoc, SourceLocation DirLoc,
+ SourceLocation EndLoc,
+ ArrayRef<const OpenACCClause *> Clauses) {
+ return new (Ctx, DC,
+ additionalSizeToAlloc<const OpenACCClause *>(Clauses.size()))
+ OpenACCDeclareDecl(DC, StartLoc, DirLoc, EndLoc, Clauses);
+}
+
+OpenACCDeclareDecl *
+OpenACCDeclareDecl::CreateDeserialized(ASTContext &Ctx, GlobalDeclID ID,
+ unsigned NumClauses) {
+ return new (Ctx, ID, additionalSizeToAlloc<const OpenACCClause *>(NumClauses))
+ OpenACCDeclareDecl(NumClauses);
+}
diff --git a/clang/lib/AST/DeclPrinter.cpp b/clang/lib/AST/DeclPrinter.cpp
index 0d51fdbc7e126..5c299567ce567 100644
--- a/clang/lib/AST/DeclPrinter.cpp
+++ b/clang/lib/AST/DeclPrinter.cpp
@@ -112,6 +112,8 @@ namespace {
void VisitNonTypeTemplateParmDecl(const NonTypeTemplateParmDecl *NTTP);
void VisitHLSLBufferDecl(HLSLBufferDecl *D);
+ void VisitOpenACCDeclareDecl(OpenACCDeclareDecl *D);
+
void printTemplateParameters(const TemplateParameterList *Params,
bool OmitTemplateKW = false);
void printTemplateArguments(llvm::ArrayRef<TemplateArgument> Args,
@@ -495,6 +497,8 @@ void DeclPrinter::VisitDeclContext(DeclContext *DC, bool Indent) {
isa<OMPDeclareMapperDecl>(*D) || isa<OMPRequiresDecl>(*D) ||
isa<OMPAllocateDecl>(*D))
Terminator = nullptr;
+ else if (isa<OpenACCDeclareDecl>(*D))
+ Terminator = nullptr;
else if (isa<ObjCMethodDecl>(*D) && cast<ObjCMethodDecl>(*D)->hasBody())
Terminator = nullptr;
else if (auto FD = dyn_cast<FunctionDecl>(*D)) {
@@ -1910,3 +1914,11 @@ void DeclPrinter::VisitNonTypeTemplateParmDecl(
/*IncludeType=*/false);
}
}
+
+void DeclPrinter::VisitOpenACCDeclareDecl(OpenACCDeclareDecl *D) {
+ if (!D->isInvalidDecl()) {
+ Out << "#pragma acc declare ";
+ OpenACCClausePrinter Printer(Out, Policy);
+ Printer.VisitClauseList(D->clauses());
+ }
+}
diff --git a/clang/lib/AST/JSONNodeDumper.cpp b/clang/lib/AST/JSONNodeDumper.cpp
index 27fd214dcee3b..d692033830191 100644
--- a/clang/lib/AST/JSONNodeDumper.cpp
+++ b/clang/lib/AST/JSONNodeDumper.cpp
@@ -1362,6 +1362,8 @@ void JSONNodeDumper::VisitSYCLUniqueStableNameExpr(
void JSONNodeDumper::VisitOpenACCAsteriskSizeExpr(
const OpenACCAsteriskSizeExpr *E) {}
+void JSONNodeDumper::VisitOpenACCDeclareDecl(const OpenACCDeclareDecl *D) {}
+
void JSONNodeDumper::VisitPredefinedExpr(const PredefinedExpr *PE) {
JOS.attribute("name", PredefinedExpr::getIdentKindName(PE->getIdentKind()));
}
diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp
index aa14ab902ba66..cef241eface91 100644
--- a/clang/lib/AST/OpenACCClause.cpp
+++ b/clang/lib/AST/OpenACCClause.cpp
@@ -40,6 +40,8 @@ bool OpenACCClauseWithVarList::classof(const OpenACCClause *C) {
OpenACCCopyInClause::classof(C) || OpenACCCopyOutClause::classof(C) ||
OpenACCReductionClause::classof(C) ||
OpenACCCreateClause::classof(C) || OpenACCDeviceClause::classof(C) ||
+ OpenACCLinkClause::classof(C) ||
+ OpenACCDeviceResidentClause::classof(C) ||
OpenACCHostClause::classof(C);
}
bool OpenACCClauseWithCondition::classof(const OpenACCClause *C) {
@@ -438,6 +440,25 @@ OpenACCCopyClause::Create(const ASTContext &C, OpenACCClauseKind Spelling,
OpenACCCopyClause(Spelling, BeginLoc, LParenLoc, VarList, EndLoc);
}
+OpenACCLinkClause *OpenACCLinkClause::Create(const ASTContext &C,
+ SourceLocation BeginLoc,
+ SourceLocation LParenLoc,
+ ArrayRef<Expr *> VarList,
+ SourceLocation EndLoc) {
+ void *Mem =
+ C.Allocate(OpenACCLinkClause::totalSizeToAlloc<Expr *>(VarList.size()));
+ return new (Mem) OpenACCLinkClause(BeginLoc, LParenLoc, VarList, EndLoc);
+}
+
+OpenACCDeviceResidentClause *OpenACCDeviceResidentClause::Create(
+ const ASTContext &C, SourceLocation BeginLoc, SourceLocation LParenLoc,
+ ArrayRef<Expr *> VarList, SourceLocation EndLoc) {
+ void *Mem = C.Allocate(
+ OpenACCDeviceResidentClause::totalSizeToAlloc<Expr *>(VarList.size()));
+ return new (Mem)
+ OpenACCDeviceResidentClause(BeginLoc, LParenLoc, VarList, EndLoc);
+}
+
OpenACCCopyInClause *
OpenACCCopyInClause::Create(const ASTContext &C, OpenACCClauseKind Spelling,
SourceLocation BeginLoc, SourceLocation LParenLoc,
@@ -754,6 +775,21 @@ void OpenACCClausePrinter::VisitCopyClause(const OpenACCCopyClause &C) {
OS << ")";
}
+void OpenACCClausePrinter::VisitLinkClause(const OpenACCLinkClause &C) {
+ OS << "link(";
+ llvm::interleaveComma(C.getVarList(), OS,
+ [&](const Expr *E) { printExpr(E); });
+ OS << ")";
+}
+
+void OpenACCClausePrinter::VisitDeviceResidentClause(
+ const OpenACCDeviceResidentClause &C) {
+ OS << "device_resident(";
+ llvm::interleaveComma(C.getVarList(), OS,
+ [&](const Expr *E) { printExpr(E); });
+ OS << ")";
+}
+
void OpenACCClausePrinter::VisitCopyInClause(const OpenACCCopyInClause &C) {
OS << C.getClauseKind() << '(';
if (C.isReadOnly())
diff --git a/clang/lib/AST/StmtPrinter.cpp b/clang/lib/AST/StmtPrinter.cpp
index c8ea7b52a6241..dac9a95e10f3d 100644
--- a/clang/lib/AST/StmtPrinter.cpp
+++ b/clang/lib/AST/StmtPrinter.cpp
@@ -17,6 +17,7 @@
#include "clang/AST/DeclBase.h"
#include "clang/AST/DeclCXX.h"
#include "clang/AST/DeclObjC.h"
+#include "clang/AST/DeclOpenACC.h"
#include "clang/AST/DeclOpenMP.h"
#include "clang/AST/DeclTemplate.h"
#include "clang/AST/Expr.h"
@@ -262,7 +263,10 @@ void StmtPrinter::VisitNullStmt(NullStmt *Node) {
void StmtPrinter::VisitDeclStmt(DeclStmt *Node) {
Indent();
PrintRawDeclStmt(Node);
- OS << ";" << NL;
+ // Certain pragma declarations shouldn't have a semi-colon after them.
+ if (!Node->isSingleDecl() || !isa<OpenACCDeclareDecl>(Node->getSingleDecl()))
+ OS << ";";
+ OS << NL;
}
void StmtPrinter::VisitCompoundStmt(CompoundStmt *Node) {
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index 2603df25ba2a4..e283a9ad4a567 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -2550,6 +2550,16 @@ void OpenACCClauseProfiler::VisitIfClause(const OpenACCIfClause &Clause) {
void OpenACCClauseProfiler::VisitCopyClause(const OpenACCCopyClause &Clause) {
VisitClauseWithVarList(Clause);
}
+
+void OpenACCClauseProfiler::VisitLinkClause(const OpenACCLinkClause &Clause) {
+ VisitClauseWithVarList(Clause);
+}
+
+void OpenACCClauseProfiler::VisitDeviceResidentClause(
+ const OpenACCDeviceResidentClause &Clause) {
+ VisitClauseWithVarList(Clause);
+}
+
void OpenACCClauseProfiler::VisitCopyInClause(
const OpenACCCopyInClause &Clause) {
VisitClauseWithVarList(Clause);
diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp
index fd1eaab9621dd..063e718454d46 100644
--- a/clang/lib/AST/TextNodeDumper.cpp
+++ b/clang/lib/AST/TextNodeDumper.cpp
@@ -417,9 +417,11 @@ void TextNodeDumper::Visit(const OpenACCClause *C) {
case OpenACCClauseKind::Device:
case OpenACCClauseKind::DeviceNum:
case OpenACCClauseKind::DefaultAsync:
+ case OpenACCClauseKind::DeviceResident:
case OpenACCClauseKind::DevicePtr:
case OpenACCClauseKind::Finalize:
case OpenACCClauseKind::FirstPrivate:
+ case OpenACCClauseKind::Link:
case OpenACCClauseKind::NoCreate:
case OpenACCClauseKind::NumGangs:
case OpenACCClauseKind::NumWorkers:
@@ -3061,6 +3063,17 @@ void TextNodeDumper::VisitOpenACCAtomicConstruct(
OS << ' ' << S->getAtomicKind();
}
+void TextNodeDumper::VisitOpenACCDeclareDecl(const OpenACCDeclareDecl *D) {
+ OS << " " << D->getDirectiveKind();
+
+ for (const OpenACCClause *C : D->clauses())
+ AddChild([=] {
+ Visit(C);
+ for (const Stmt *S : C->children())
+ AddChild([=] { Visit(S); });
+ });
+}
+
void TextNodeDumper::VisitEmbedExpr(const EmbedExpr *S) {
AddChild("begin", [=] { OS << S->getStartingElementPos(); });
AddChild("number of elements", [=] { OS << S->getDataElementCount(); });
diff --git a/clang/lib/CodeGen/CGDecl.cpp b/clang/lib/CodeGen/CGDecl.cpp
index 668282a6ab1a8..9cd5885aaae51 100644
--- a/clang/lib/CodeGen/CGDecl.cpp
+++ b/clang/lib/CodeGen/CGDecl.cpp
@@ -27,6 +27,7 @@
#include "clang/AST/CharUnits.h"
#include "clang/AST/Decl.h"
#include "clang/AST/DeclObjC.h"
+#include "clang/AST/DeclOpenACC.h"
#include "clang/AST/DeclOpenMP.h"
#include "clang/Basic/CodeGenOptions.h"
#include "clang/Basic/TargetInfo.h"
@@ -177,6 +178,9 @@ void CodeGenFunction::EmitDecl(const Decl &D) {
case Decl::OMPDeclareMapper:
return CGM.EmitOMPDeclareMapper(cast<OMPDeclareMapperDecl>(&D), this);
+ case Decl::OpenACCDeclare:
+ return CGM.EmitOpenACCDeclare(cast<OpenACCDeclareDecl>(&D), this);
+
case Decl::Typedef: // typedef int X;
case Decl::TypeAlias: { // using X = int; [C++0x]
QualType Ty = cast<TypedefNameDecl>(D).getUnderlyingType();
@@ -2843,6 +2847,11 @@ void CodeGenModule::EmitOMPDeclareMapper(const OMPDeclareMapperDecl *D,
getOpenMPRuntime().emitUserDefinedMapper(D, CGF);
}
+void CodeGenModule::EmitOpenACCDeclare(const OpenACCDeclareDecl *D,
+ CodeGenFunction *CGF) {
+ // This is a no-op, we cna just ignore these declarations.
+}
+
void CodeGenModule::EmitOMPRequiresDecl(const OMPRequiresDecl *D) {
getOpenMPRuntime().processRequiresDirective(D);
}
diff --git a/clang/lib/CodeGen/CodeGenModule.h b/clang/lib/CodeGen/CodeGenModule.h
index 4a269f622ece4..5726d7939b3c3 100644
--- a/clang/lib/CodeGen/CodeGenModule.h
+++ b/clang/lib/CodeGen/CodeGenModule.h
@@ -1569,6 +1569,9 @@ class CodeGenModule : public CodeGenTypeCache {
void EmitOMPDeclareMapper(const OMPDeclareMapperDecl *D,
CodeGenFunction *CGF = nullptr);
+ // Emit code for the OpenACC Declare declaration.
+ void EmitOpenACCDeclare(const OpenACCDeclareDecl *D, CodeGenFunction *CGF);
+
/// Emit a code for requires directive.
/// \param D Requires declaration
void EmitOMPRequiresDecl(const OMPRequiresDecl *D);
diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
index d036971d2fc31..5eefd3ffb092e 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -981,7 +981,7 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
case OpenACCClauseKind::PresentOrCopyIn: {
bool IsReadOnly = tryParseAndConsumeSpecialTokenKind(
*this, OpenACCSpecialTokenKind::ReadOnly, ClauseKind);
- ParsedClause.setVarListDetails(ParseOpenACCVarList(ClauseKind),
+ ParsedClause.setVarListDetails(ParseOpenACCVarList(DirKind, ClauseKind),
IsReadOnly,
/*IsZero=*/false);
break;
@@ -994,7 +994,7 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
case OpenACCClauseKind::PresentOrCopyOut: {
bool IsZero = tryParseAndConsumeSpecialTokenKind(
*this, OpenACCSpecialTokenKind::Zero, ClauseKind);
- ParsedClause.setVarListDetails(ParseOpenACCVarList(ClauseKind),
+ ParsedClause.setVarListDetails(ParseOpenACCVarList(DirKind, ClauseKind),
/*IsReadOnly=*/false, IsZero);
break;
}
@@ -1002,7 +1002,8 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
// If we're missing a clause-kind (or it is invalid), see if we can parse
// the var-list anyway.
OpenACCReductionOperator Op = ParseReductionOperator(*this);
- ParsedClause.setReductionDetails(Op, ParseOpenACCVarList(ClauseKind));
+ ParsedClause.setReductionDetails(
+ Op, ParseOpenACCVarList(DirKind, ClauseKind));
break;
}
case OpenACCClauseKind::Self:
@@ -1013,21 +1014,13 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
[[fallthrough]];
case OpenACCClauseKind::Device:
case OpenACCClauseKind::Host:
- ParsedClause.setVarListDetails(ParseOpenACCVarList(ClauseKind),
- /*IsReadOnly=*/false, /*IsZero=*/false);
- break;
case OpenACCClauseKind::DeviceResident:
case OpenACCClauseKind::Link:
- ParseOpenACCVarList(ClauseKind);
- break;
case OpenACCClauseKind::Attach:
case OpenACCClauseKind::Delete:
case OpenACCClauseKind::Detach:
case OpenACCClauseKind::DevicePtr:
case OpenACCClauseKind::UseDevice:
- ParsedClause.setVarListDetails(ParseOpenACCVarList(ClauseKind),
- /*IsReadOnly=*/false, /*IsZero=*/false);
- break;
case OpenACCClauseKind::Copy:
case OpenACCClauseKind::PCopy:
case OpenACCClauseKind::PresentOrCopy:
@@ -1035,7 +1028,7 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
case OpenACCClauseKind::NoCreate:
case OpenACCClauseKind::Present:
case OpenACCClauseKind::Private:
- ParsedClause.setVarListDetails(ParseOpenACCVarList(ClauseKind),
+ ParsedClause.setVarListDetails(ParseOpenACCVarList(DirKind, ClauseKind),
/*IsReadOnly=*/false, /*IsZero=*/false);
break;
case OpenACCClauseKind::Collapse: {
@@ -1362,7 +1355,8 @@ ExprResult Parser::ParseOpenACCBindClauseArgument() {
/// - an array element
/// - a member of a composite variable
/// - a common block name between slashes (fortran only)
-Parser::OpenACCVarParseResult Parser::ParseOpenACCVar(OpenACCClauseKind CK) {
+Parser::OpenACCVarParseResult Parser::ParseOpenACCVar(OpenACCDirectiveKind DK,
+ OpenACCClauseKind CK) {
OpenACCArraySectionRAII ArraySections(*this);
ExprResult Res = ParseAssignmentExpression();
@@ -1373,15 +1367,16 @@ Parser::OpenACCVarParseResult Parser::ParseOpenACCVar(OpenACCClauseKind CK) {
if (!Res.isUsable())
return {Res, OpenACCParseCanContinue::Can};
- Res = getActions().OpenACC().ActOnVar(CK, Res.get());
+ Res = getActions().OpenACC().ActOnVar(DK, CK, Res.get());
return {Res, OpenACCParseCanContinue::Can};
}
-llvm::SmallVector<Expr *> Parser::ParseOpenACCVarList(OpenACCClauseKind CK) {
+llvm::SmallVector<Expr *> Parser::ParseOpenACCVarList(OpenACCDirectiveKind DK,
+ OpenACCClauseKind CK) {
llvm::SmallVector<Expr *> Vars;
- auto [Res, CanContinue] = ParseOpenACCVar(CK);
+ auto [Res, CanContinue] = ParseOpenACCVar(DK, CK);
if (Res.isUsable()) {
Vars.push_back(Res.get());
} else if (CanContinue == OpenACCParseCanContinue::Cannot) {
@@ -1392,7 +1387,7 @@ llvm::SmallVector<Expr *> Parser::ParseOpenACCVarList(OpenACCClauseKind CK) {
while (!getCurToken().isOneOf(tok::r_paren, tok::annot_pragma_openacc_end)) {
ExpectAndConsume(tok::comma);
- auto [Res, CanContinue] = ParseOpenACCVar(CK);
+ auto [Res, CanContinue] = ParseOpenACCVar(DK, CK);
if (Res.isUsable()) {
Vars.push_back(Res.get());
@@ -1426,7 +1421,7 @@ void Parser::ParseOpenACCCacheVarList() {
// ParseOpenACCVarList should leave us before a r-paren, so no need to skip
// anything here.
- ParseOpenACCVarList(OpenACCClauseKind::Invalid);
+ ParseOpenACCVarList(OpenACCDirectiveKind::Cache, OpenACCClauseKind::Invalid);
}
Parser::OpenACCDirectiveParseInfo
@@ -1523,8 +1518,9 @@ Parser::DeclGroupPtrTy Parser::ParseOpenACCDirectiveDecl() {
DirInfo.StartLoc))
return nullptr;
- // TODO OpenACC: Do whatever decl parsing is required here.
- return DeclGroupPtrTy::make(getActions().OpenACC().ActOnEndDeclDirective());
+ return DeclGroupPtrTy::make(getActions().OpenACC().ActOnEndDeclDirective(
+ DirInfo.DirKind, DirInfo.StartLoc, DirInfo.DirLoc, DirInfo.EndLoc,
+ DirInfo.Clauses));
}
// Parse OpenACC Directive on a Statement.
diff --git a/clang/lib/Sema/SemaBase.cpp b/clang/lib/Sema/SemaBase.cpp
index 5c24f21b469b0..85c4a0ab40fed 100644
--- a/clang/lib/Sema/SemaBase.cpp
+++ b/clang/lib/Sema/SemaBase.cpp
@@ -9,6 +9,7 @@ SemaBase::SemaBase(Sema &S) : SemaRef(S) {}
ASTContext &SemaBase::getASTContext() const { return SemaRef.Context; }
DiagnosticsEngine &SemaBase::getDiagnostics() const { return SemaRef.Diags; }
const LangOptions &SemaBase::getLangOpts() const { return SemaRef.LangOpts; }
+DeclContext *SemaBase::getCurContext() const { return SemaRef.CurContext; }
SemaBase::ImmediateDiagBuilder::~ImmediateDiagBuilder() {
// If we aren't active, there is nothing to do.
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 2d2f8ddf4652b..59583e73952cf 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -11,11 +11,12 @@
///
//===----------------------------------------------------------------------===//
+#include "clang/Sema/SemaOpenACC.h"
+#include "clang/AST/DeclOpenACC.h"
#include "clang/AST/StmtOpenACC.h"
#include "clang/Basic/DiagnosticSema.h"
#include "clang/Basic/OpenACCKinds.h"
#include "clang/Sema/Sema.h"
-#include "clang/Sema/SemaOpenACC.h"
#include "llvm/ADT/StringExtras.h"
#include "llvm/Support/Casting.h"
@@ -334,6 +335,7 @@ void SemaOpenACC::ActOnConstruct(OpenACCDirectiveKind K,
case OpenACCDirectiveKind::Set:
case OpenACCDirectiveKind::Update:
case OpenACCDirectiveKind::Atomic:
+ case OpenACCDirectiveKind::Declare:
// Nothing to do here, there is no real legalization that needs to happen
// here as these constructs do not take any arguments.
break;
@@ -478,15 +480,20 @@ bool SemaOpenACC::CheckVarIsPointerType(OpenACCClauseKind ClauseKind,
return false;
}
-ExprResult SemaOpenACC::ActOnVar(OpenACCClauseKind CK, Expr *VarExpr) {
+ExprResult SemaOpenACC::ActOnVar(OpenACCDirectiveKind DK, OpenACCClauseKind CK,
+ Expr *VarExpr) {
Expr *CurVarExpr = VarExpr->IgnoreParenImpCasts();
// 'use_device' doesn't allow array subscript or array sections.
// OpenACC3.3 2.8:
// A 'var' in a 'use_device' clause must be the name of a variable or array.
- if (CK == OpenACCClauseKind::UseDevice &&
+ // OpenACC3.3 2.13:
+ // A 'var' in a 'declare' directive must be a variable or array name.
+ if ((CK == OpenACCClauseKind::UseDevice ||
+ DK == OpenACCDirectiveKind::Declare) &&
isa<ArraySectionExpr, ArraySubscriptExpr>(CurVarExpr)) {
- Diag(VarExpr->getExprLoc(), diag::err_acc_not_a_var_ref_use_device);
+ Diag(VarExpr->getExprLoc(), diag::err_acc_not_a_var_ref_use_device_declare)
+ << (DK == OpenACCDirectiveKind::Declare);
return ExprError();
}
@@ -510,20 +517,30 @@ ExprResult SemaOpenACC::ActOnVar(OpenACCClauseKind CK, Expr *VarExpr) {
// If CK is a Reduction, this special cases for OpenACC3.3 2.5.15: "A var in a
// reduction clause must be a scalar variable name, an aggregate variable
// name, an array element, or a subarray.
- // If CK is a 'use_device', this also isn't valid, as it isn' the name of a
- // variable or array.
+ // If CK is a 'use_device', this also isn't valid, as it isn't the name of a
+ // variable or array, if not done as a member expr.
// A MemberExpr that references a Field is valid for other clauses.
- if (CK != OpenACCClauseKind::Reduction &&
- CK != OpenACCClauseKind::UseDevice) {
- if (const auto *ME = dyn_cast<MemberExpr>(CurVarExpr)) {
- if (isa<FieldDecl>(ME->getMemberDecl()->getCanonicalDecl()))
+ if (const auto *ME = dyn_cast<MemberExpr>(CurVarExpr)) {
+ if (isa<FieldDecl>(ME->getMemberDecl()->getCanonicalDecl())) {
+ if (DK == OpenACCDirectiveKind::Declare ||
+ CK == OpenACCClauseKind::Reduction ||
+ CK == OpenACCClauseKind::UseDevice) {
+
+ // We can allow 'member expr' if the 'this' is implicit in the case of
+ // declare, reduction, and use_device.
+ const auto *This = dyn_cast<CXXThisExpr>(ME->getBase());
+ if (This && This->isImplicit())
+ return VarExpr;
+ } else {
return VarExpr;
+ }
}
}
- // Referring to 'this' is ok for the most part, but for 'use_device' doesn't
- // fall into 'variable or array name'
- if (CK != OpenACCClauseKind::UseDevice && isa<CXXThisExpr>(CurVarExpr))
+ // Referring to 'this' is ok for the most part, but for 'use_device'/'declare'
+ // doesn't fall into 'variable or array name'
+ if (CK != OpenACCClauseKind::UseDevice &&
+ DK != OpenACCDirectiveKind::Declare && isa<CXXThisExpr>(CurVarExpr))
return VarExpr;
// Nothing really we can do here, as these are dependent. So just return they
@@ -538,8 +555,12 @@ ExprResult SemaOpenACC::ActOnVar(OpenACCClauseKind CK, Expr *VarExpr) {
if (isa<RecoveryExpr>(CurVarExpr))
return ExprError();
- if (CK == OpenACCClauseKind::UseDevice)
- Diag(VarExpr->getExprLoc(), diag::err_acc_not_a_var_ref_use_device);
+ if (DK == OpenACCDirectiveKind::Declare)
+ Diag(VarExpr->getExprLoc(), diag::err_acc_not_a_var_ref_use_device_declare)
+ << /*declare*/ 1;
+ else if (CK == OpenACCClauseKind::UseDevice)
+ Diag(VarExpr->getExprLoc(), diag::err_acc_not_a_var_ref_use_device_declare)
+ << /*use_device*/ 0;
else
Diag(VarExpr->getExprLoc(), diag::err_acc_not_a_var_ref)
<< (CK != OpenACCClauseKind::Reduction);
@@ -1413,6 +1434,12 @@ std::string GetListOfClauses(llvm::ArrayRef<OpenACCClauseKind> Clauses) {
bool SemaOpenACC::ActOnStartStmtDirective(
OpenACCDirectiveKind K, SourceLocation StartLoc,
ArrayRef<const OpenACCClause *> Clauses) {
+
+ // Declaration directives an appear in a statement location, so call into that
+ // function here.
+ if (K == OpenACCDirectiveKind::Declare || K == OpenACCDirectiveKind::Routine)
+ return ActOnStartDeclDirective(K, StartLoc);
+
SemaRef.DiscardCleanupsInEvaluationContext();
SemaRef.PopExpressionEvaluationContext();
@@ -1597,6 +1624,14 @@ StmtResult SemaOpenACC::ActOnEndStmtDirective(
getASTContext(), StartLoc, DirLoc, AtomicKind, EndLoc,
AssocStmt.isUsable() ? AssocStmt.get() : nullptr);
}
+ 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.
+
+ DeclGroupRef DR =
+ ActOnEndDeclDirective(K, StartLoc, DirLoc, EndLoc, Clauses);
+ return SemaRef.ActOnDeclStmt(DeclGroupPtrTy::make(DR), StartLoc, EndLoc);
+ }
}
llvm_unreachable("Unhandled case in directive handling?");
}
@@ -1683,7 +1718,33 @@ bool SemaOpenACC::ActOnStartDeclDirective(OpenACCDirectiveKind K,
return diagnoseConstructAppertainment(*this, K, StartLoc, /*IsStmt=*/false);
}
-DeclGroupRef SemaOpenACC::ActOnEndDeclDirective() { return DeclGroupRef{}; }
+DeclGroupRef SemaOpenACC::ActOnEndDeclDirective(
+ OpenACCDirectiveKind K, SourceLocation StartLoc, SourceLocation DirLoc,
+ SourceLocation EndLoc, ArrayRef<OpenACCClause *> Clauses) {
+ switch (K) {
+ default:
+ case OpenACCDirectiveKind::Invalid:
+ return DeclGroupRef{};
+ case OpenACCDirectiveKind::Declare: {
+ // OpenACC3.3 2.13: At least one clause must appear on a declare directive.
+ if (Clauses.empty()) {
+ Diag(EndLoc, diag::err_acc_declare_required_clauses);
+ // No reason to add this to the AST, as we would just end up trying to
+ // instantiate this, which would double-diagnose here, which we wouldn't
+ // want to do.
+ return DeclGroupRef{};
+ }
+
+ auto *DeclareDecl = OpenACCDeclareDecl::Create(
+ getASTContext(), getCurContext(), StartLoc, DirLoc, EndLoc, Clauses);
+ DeclareDecl->setAccess(AS_public);
+ getCurContext()->addDecl(DeclareDecl);
+ return DeclGroupRef{DeclareDecl};
+ }
+ }
+
+ llvm_unreachable("unhandled case in directive handling?");
+}
ExprResult
SemaOpenACC::BuildOpenACCAsteriskSizeExpr(SourceLocation AsteriskLoc) {
diff --git a/clang/lib/Sema/SemaOpenACCClause.cpp b/clang/lib/Sema/SemaOpenACCClause.cpp
index 1e74f126c31ce..0805779d74bdc 100644
--- a/clang/lib/Sema/SemaOpenACCClause.cpp
+++ b/clang/lib/Sema/SemaOpenACCClause.cpp
@@ -10,8 +10,9 @@
///
//===----------------------------------------------------------------------===//
-#include "clang/AST/OpenACCClause.h"
#include "clang/AST/DeclCXX.h"
+#include "clang/AST/ExprCXX.h"
+#include "clang/AST/OpenACCClause.h"
#include "clang/Basic/DiagnosticSema.h"
#include "clang/Basic/OpenACCKinds.h"
#include "clang/Sema/SemaOpenACC.h"
@@ -192,6 +193,7 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
case OpenACCDirectiveKind::Kernels:
case OpenACCDirectiveKind::Data:
case OpenACCDirectiveKind::EnterData:
+ case OpenACCDirectiveKind::Declare:
case OpenACCDirectiveKind::ParallelLoop:
case OpenACCDirectiveKind::SerialLoop:
case OpenACCDirectiveKind::KernelsLoop:
@@ -424,6 +426,22 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
return false;
}
}
+ case OpenACCClauseKind::Link: {
+ switch (DirectiveKind) {
+ case OpenACCDirectiveKind::Declare:
+ return true;
+ default:
+ return false;
+ }
+ }
+ case OpenACCClauseKind::DeviceResident: {
+ switch (DirectiveKind) {
+ case OpenACCDirectiveKind::Declare:
+ return true;
+ default:
+ return false;
+ }
+ }
case OpenACCClauseKind::UseDevice: {
switch (DirectiveKind) {
@@ -588,8 +606,17 @@ bool checkValidAfterDeviceType(
// with the one being currently implemented/only updated after the entire
// construct has been implemented.
bool isDirectiveKindImplemented(OpenACCDirectiveKind DK) {
- return DK != OpenACCDirectiveKind::Declare &&
- DK != OpenACCDirectiveKind::Routine;
+ return DK != OpenACCDirectiveKind::Routine;
+}
+
+// GCC looks through linkage specs, but not the other transparent declaration
+// contexts for 'declare' restrictions, so this helper function helps get us
+// through that.
+const DeclContext *removeLinkageSpecDC(const DeclContext *DC) {
+ while (isa<LinkageSpecDecl>(DC))
+ DC = DC->getParent();
+
+ return DC;
}
class SemaOpenACCClauseVisitor {
@@ -1006,16 +1033,15 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitNoCreateClause(
OpenACCClause *SemaOpenACCClauseVisitor::VisitPresentClause(
SemaOpenACC::OpenACCParsedClause &Clause) {
- // Restrictions only properly implemented on 'compute'/'combined'/'data'
- // constructs, and 'compute'/'combined'/'data' constructs are the only
- // construct that can do anything with this yet, so skip/treat as
- // unimplemented in this case.
- if (!isDirectiveKindImplemented(Clause.getDirectiveKind()))
- return isNotImplemented();
// ActOnVar ensured that everything is a valid variable reference, so there
// really isn't anything to do here. GCC does some duplicate-finding, though
// it isn't apparent in the standard where this is justified.
+ // 'declare' has some restrictions that need to be enforced separately, so
+ // check it here.
+ if (SemaRef.CheckDeclareClause(Clause))
+ return nullptr;
+
return OpenACCPresentClause::Create(Ctx, Clause.getBeginLoc(),
Clause.getLParenLoc(),
Clause.getVarList(), Clause.getEndLoc());
@@ -1045,33 +1071,58 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitDeviceClause(
OpenACCClause *SemaOpenACCClauseVisitor::VisitCopyClause(
SemaOpenACC::OpenACCParsedClause &Clause) {
- // Restrictions only properly implemented on 'compute'/'combined'/'data'
- // constructs, and 'compute'/'combined'/'data' constructs are the only
- // construct that can do anything with this yet, so skip/treat as
- // unimplemented in this case.
- if (!isDirectiveKindImplemented(Clause.getDirectiveKind()))
- return isNotImplemented();
// ActOnVar ensured that everything is a valid variable reference, so there
// really isn't anything to do here. GCC does some duplicate-finding, though
// it isn't apparent in the standard where this is justified.
+ // 'declare' has some restrictions that need to be enforced separately, so
+ // check it here.
+ if (SemaRef.CheckDeclareClause(Clause))
+ return nullptr;
+
return OpenACCCopyClause::Create(
Ctx, Clause.getClauseKind(), Clause.getBeginLoc(), Clause.getLParenLoc(),
Clause.getVarList(), Clause.getEndLoc());
}
+OpenACCClause *SemaOpenACCClauseVisitor::VisitLinkClause(
+ SemaOpenACC::OpenACCParsedClause &Clause) {
+ // 'declare' has some restrictions that need to be enforced separately, so
+ // check it here.
+ if (SemaRef.CheckDeclareClause(Clause))
+ return nullptr;
+
+ Clause.setVarListDetails(SemaRef.CheckLinkClauseVarList(Clause.getVarList()),
+ /*IsReadOnly=*/false, /*IsZero=*/false);
+
+ return OpenACCLinkClause::Create(Ctx, Clause.getBeginLoc(),
+ Clause.getLParenLoc(), Clause.getVarList(),
+ Clause.getEndLoc());
+}
+
+OpenACCClause *SemaOpenACCClauseVisitor::VisitDeviceResidentClause(
+ SemaOpenACC::OpenACCParsedClause &Clause) {
+ // 'declare' has some restrictions that need to be enforced separately, so
+ // check it here.
+ if (SemaRef.CheckDeclareClause(Clause))
+ return nullptr;
+
+ return OpenACCDeviceResidentClause::Create(
+ Ctx, Clause.getBeginLoc(), Clause.getLParenLoc(), Clause.getVarList(),
+ Clause.getEndLoc());
+}
+
OpenACCClause *SemaOpenACCClauseVisitor::VisitCopyInClause(
SemaOpenACC::OpenACCParsedClause &Clause) {
- // Restrictions only properly implemented on 'compute'/'combined'/'data'
- // constructs, and 'compute'/'combined'/'data' constructs are the only
- // construct that can do anything with this yet, so skip/treat as
- // unimplemented in this case.
- if (!isDirectiveKindImplemented(Clause.getDirectiveKind()))
- return isNotImplemented();
// ActOnVar ensured that everything is a valid variable reference, so there
// really isn't anything to do here. GCC does some duplicate-finding, though
// it isn't apparent in the standard where this is justified.
+ // 'declare' has some restrictions that need to be enforced separately, so
+ // check it here.
+ if (SemaRef.CheckDeclareClause(Clause))
+ return nullptr;
+
return OpenACCCopyInClause::Create(
Ctx, Clause.getClauseKind(), Clause.getBeginLoc(), Clause.getLParenLoc(),
Clause.isReadOnly(), Clause.getVarList(), Clause.getEndLoc());
@@ -1079,16 +1130,15 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitCopyInClause(
OpenACCClause *SemaOpenACCClauseVisitor::VisitCopyOutClause(
SemaOpenACC::OpenACCParsedClause &Clause) {
- // Restrictions only properly implemented on 'compute'/'combined'/'data'
- // constructs, and 'compute'/'combined'/'data' constructs are the only
- // construct that can do anything with this yet, so skip/treat as
- // unimplemented in this case.
- if (!isDirectiveKindImplemented(Clause.getDirectiveKind()))
- return isNotImplemented();
// ActOnVar ensured that everything is a valid variable reference, so there
// really isn't anything to do here. GCC does some duplicate-finding, though
// it isn't apparent in the standard where this is justified.
+ // 'declare' has some restrictions that need to be enforced separately, so
+ // check it here.
+ if (SemaRef.CheckDeclareClause(Clause))
+ return nullptr;
+
return OpenACCCopyOutClause::Create(
Ctx, Clause.getClauseKind(), Clause.getBeginLoc(), Clause.getLParenLoc(),
Clause.isZero(), Clause.getVarList(), Clause.getEndLoc());
@@ -1100,6 +1150,11 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitCreateClause(
// really isn't anything to do here. GCC does some duplicate-finding, though
// it isn't apparent in the standard where this is justified.
+ // 'declare' has some restrictions that need to be enforced separately, so
+ // check it here.
+ if (SemaRef.CheckDeclareClause(Clause))
+ return nullptr;
+
return OpenACCCreateClause::Create(
Ctx, Clause.getClauseKind(), Clause.getBeginLoc(), Clause.getLParenLoc(),
Clause.isZero(), Clause.getVarList(), Clause.getEndLoc());
@@ -1156,13 +1211,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitUseDeviceClause(
OpenACCClause *SemaOpenACCClauseVisitor::VisitDevicePtrClause(
SemaOpenACC::OpenACCParsedClause &Clause) {
- // Restrictions only properly implemented on 'compute'/'combined'/'data'
- // constructs, and 'compute'/'combined'/'data' constructs are the only
- // construct that can do anything with this yet, so skip/treat as
- // unimplemented in this case.
- if (!isDirectiveKindImplemented(Clause.getDirectiveKind()))
- return isNotImplemented();
-
// ActOnVar ensured that everything is a valid variable reference, but we
// still have to make sure it is a pointer type.
llvm::SmallVector<Expr *> VarList{Clause.getVarList()};
@@ -1172,6 +1220,11 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitDevicePtrClause(
Clause.setVarListDetails(VarList,
/*IsReadOnly=*/false, /*IsZero=*/false);
+ // 'declare' has some restrictions that need to be enforced separately, so
+ // check it here.
+ if (SemaRef.CheckDeclareClause(Clause))
+ return nullptr;
+
return OpenACCDevicePtrClause::Create(
Ctx, Clause.getBeginLoc(), Clause.getLParenLoc(), Clause.getVarList(),
Clause.getEndLoc());
@@ -2243,3 +2296,136 @@ OpenACCClause *SemaOpenACC::CheckReductionClause(
getASTContext(), BeginLoc, LParenLoc, ReductionOp, Vars, EndLoc);
return Ret;
}
+
+llvm::SmallVector<Expr *>
+SemaOpenACC::CheckLinkClauseVarList(ArrayRef<Expr *> VarExprs) {
+ const DeclContext *DC = removeLinkageSpecDC(getCurContext());
+
+ // Link has no special restrictions on its var list unless it is not at NS/TU
+ // scope.
+ if (isa<NamespaceDecl, TranslationUnitDecl>(DC))
+ return llvm::SmallVector<Expr *>(VarExprs);
+
+ llvm::SmallVector<Expr *> NewVarList;
+
+ for (Expr *VarExpr : VarExprs) {
+ if (isa<DependentScopeDeclRefExpr, CXXDependentScopeMemberExpr>(VarExpr)) {
+ NewVarList.push_back(VarExpr);
+ continue;
+ }
+
+ // Field decls can't be global, nor extern, and declare can't refer to
+ // non-static fields in class-scope, so this always fails the scope check.
+ // BUT for now we add this so it gets diagnosed by the general 'declare'
+ // rules.
+ if (isa<MemberExpr>(VarExpr)) {
+ NewVarList.push_back(VarExpr);
+ continue;
+ }
+
+ const auto *DRE = cast<DeclRefExpr>(VarExpr);
+ const VarDecl *Var = dyn_cast<VarDecl>(DRE->getDecl());
+
+ if (!Var || !Var->hasExternalStorage())
+ Diag(VarExpr->getBeginLoc(), diag::err_acc_link_not_extern);
+ else
+ NewVarList.push_back(VarExpr);
+ }
+
+ return NewVarList;
+}
+bool SemaOpenACC::CheckDeclareClause(SemaOpenACC::OpenACCParsedClause &Clause) {
+
+ if (Clause.getDirectiveKind() != OpenACCDirectiveKind::Declare)
+ return false;
+
+ const DeclContext *DC = removeLinkageSpecDC(getCurContext());
+
+ // Whether this is 'create', 'copyin', 'deviceptr', 'device_resident', or
+ // 'link', which have 2 special rules.
+ bool IsSpecialClause =
+ Clause.getClauseKind() == OpenACCClauseKind::Create ||
+ Clause.getClauseKind() == OpenACCClauseKind::CopyIn ||
+ Clause.getClauseKind() == OpenACCClauseKind::DevicePtr ||
+ Clause.getClauseKind() == OpenACCClauseKind::DeviceResident ||
+ Clause.getClauseKind() == OpenACCClauseKind::Link;
+
+ // OpenACC 3.3 2.13:
+ // In C or C++ global or namespace scope, only 'create',
+ // 'copyin', 'deviceptr', 'device_resident', or 'link' clauses are
+ // allowed.
+ if (!IsSpecialClause && isa<NamespaceDecl, TranslationUnitDecl>(DC)) {
+ return Diag(Clause.getBeginLoc(), diag::err_acc_declare_clause_at_global)
+ << Clause.getClauseKind();
+ }
+
+ llvm::SmallVector<Expr *> FilteredVarList;
+ const DeclaratorDecl *CurDecl = nullptr;
+ for (Expr *VarExpr : Clause.getVarList()) {
+ if (isa<DependentScopeDeclRefExpr, CXXDependentScopeMemberExpr>(VarExpr)) {
+ // There isn't really anything we can do here, so we add them anyway and
+ // we can check them again when we instantiate this.
+ } else if (const auto *MemExpr = dyn_cast<MemberExpr>(VarExpr)) {
+ FieldDecl *FD =
+ cast<FieldDecl>(MemExpr->getMemberDecl()->getCanonicalDecl());
+ CurDecl = FD;
+
+ if (removeLinkageSpecDC(
+ FD->getLexicalDeclContext()->getPrimaryContext()) != DC) {
+ Diag(MemExpr->getBeginLoc(), diag::err_acc_declare_same_scope)
+ << Clause.getClauseKind();
+ continue;
+ }
+ } else {
+ const auto *DRE = cast<DeclRefExpr>(VarExpr);
+ const VarDecl *Var = dyn_cast<VarDecl>(DRE->getDecl());
+ if (Var)
+ CurDecl = Var->getCanonicalDecl();
+
+ // OpenACC3.3 2.13:
+ // A 'declare' directive must be in the same scope as the declaration of
+ // any var that appears in the clauses of the directive or any scope
+ // within a C/C++ function.
+ // We can't really check 'scope' here, so we check declaration context,
+ // which is a reasonable approximation, but misses scopes inside of
+ // functions.
+ if (removeLinkageSpecDC(Var->getCanonicalDecl()
+ ->getLexicalDeclContext()
+ ->getPrimaryContext()) != DC) {
+ Diag(VarExpr->getBeginLoc(), diag::err_acc_declare_same_scope)
+ << Clause.getClauseKind();
+ continue;
+ }
+ // OpenACC3.3 2.13:
+ // C and C++ extern variables may only appear in 'create',
+ // 'copyin', 'deviceptr', 'device_resident', or 'link' clauses on a
+ // 'declare' directive.
+ if (!IsSpecialClause && Var && Var->hasExternalStorage()) {
+ Diag(VarExpr->getBeginLoc(), diag::err_acc_declare_extern)
+ << Clause.getClauseKind();
+ continue;
+ }
+
+ // OpenACC3.3 2.13:
+ // A var may appear at most once in all the clauses of declare
+ // directives for a function, subroutine, program, or module.
+
+ if (CurDecl) {
+ auto Itr = DeclareVarReferences.find(CurDecl);
+ if (Itr != DeclareVarReferences.end()) {
+ Diag(VarExpr->getBeginLoc(), diag::err_acc_multiple_references)
+ << Clause.getClauseKind();
+ Diag(Itr->second, diag::note_acc_previous_reference);
+ continue;
+ } else {
+ DeclareVarReferences[CurDecl] = VarExpr->getBeginLoc();
+ }
+ }
+ }
+ FilteredVarList.push_back(VarExpr);
+ }
+
+ Clause.setVarListDetails(FilteredVarList, Clause.isReadOnly(),
+ Clause.isZero());
+ return false;
+}
diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
index 1cdf80898bfca..dd894df851488 100644
--- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
@@ -999,6 +999,234 @@ TemplateDeclInstantiator::VisitNamespaceDecl(NamespaceDecl *D) {
llvm_unreachable("Namespaces cannot be instantiated");
}
+namespace {
+class OpenACCDeclClauseInstantiator final
+ : public OpenACCClauseVisitor<OpenACCDeclClauseInstantiator> {
+ Sema &SemaRef;
+ const MultiLevelTemplateArgumentList &MLTAL;
+ SemaOpenACC::OpenACCParsedClause &ParsedClause;
+ OpenACCClause *NewClause = nullptr;
+
+public:
+ OpenACCDeclClauseInstantiator(Sema &S,
+ const MultiLevelTemplateArgumentList &MLTAL,
+ SemaOpenACC::OpenACCParsedClause &ParsedClause)
+ : SemaRef(S), MLTAL(MLTAL), ParsedClause(ParsedClause) {}
+
+ OpenACCClause *CreatedClause() { return NewClause; }
+#define VISIT_CLAUSE(CLAUSE_NAME) \
+ void Visit##CLAUSE_NAME##Clause(const OpenACC##CLAUSE_NAME##Clause &Clause);
+#include "clang/Basic/OpenACCClauses.def"
+
+ llvm::SmallVector<Expr *> VisitVarList(ArrayRef<Expr *> VarList) {
+ llvm::SmallVector<Expr *> InstantiatedVarList;
+ for (Expr *CurVar : VarList) {
+ ExprResult Res = SemaRef.SubstExpr(CurVar, MLTAL);
+
+ if (!Res.isUsable())
+ continue;
+
+ Res = SemaRef.OpenACC().ActOnVar(ParsedClause.getDirectiveKind(),
+ ParsedClause.getClauseKind(), Res.get());
+
+ if (Res.isUsable())
+ InstantiatedVarList.push_back(Res.get());
+ }
+ return InstantiatedVarList;
+ }
+};
+
+#define CLAUSE_NOT_ON_DECLS(CLAUSE_NAME) \
+ void OpenACCDeclClauseInstantiator::Visit##CLAUSE_NAME##Clause( \
+ const OpenACC##CLAUSE_NAME##Clause &) { \
+ llvm_unreachable("Clause type invalid on declaration construct, or " \
+ "instantiation not implemented"); \
+ }
+
+CLAUSE_NOT_ON_DECLS(Auto)
+CLAUSE_NOT_ON_DECLS(Async)
+CLAUSE_NOT_ON_DECLS(Attach)
+CLAUSE_NOT_ON_DECLS(Collapse)
+CLAUSE_NOT_ON_DECLS(Default)
+CLAUSE_NOT_ON_DECLS(DefaultAsync)
+CLAUSE_NOT_ON_DECLS(Delete)
+CLAUSE_NOT_ON_DECLS(Detach)
+CLAUSE_NOT_ON_DECLS(Device)
+CLAUSE_NOT_ON_DECLS(DeviceNum)
+CLAUSE_NOT_ON_DECLS(DeviceType)
+CLAUSE_NOT_ON_DECLS(Finalize)
+CLAUSE_NOT_ON_DECLS(FirstPrivate)
+CLAUSE_NOT_ON_DECLS(Gang)
+CLAUSE_NOT_ON_DECLS(Host)
+CLAUSE_NOT_ON_DECLS(If)
+CLAUSE_NOT_ON_DECLS(IfPresent)
+CLAUSE_NOT_ON_DECLS(Independent)
+CLAUSE_NOT_ON_DECLS(NoCreate)
+CLAUSE_NOT_ON_DECLS(NumGangs)
+CLAUSE_NOT_ON_DECLS(NumWorkers)
+CLAUSE_NOT_ON_DECLS(Private)
+CLAUSE_NOT_ON_DECLS(Reduction)
+CLAUSE_NOT_ON_DECLS(Self)
+CLAUSE_NOT_ON_DECLS(Seq)
+CLAUSE_NOT_ON_DECLS(Tile)
+CLAUSE_NOT_ON_DECLS(UseDevice)
+CLAUSE_NOT_ON_DECLS(Vector)
+CLAUSE_NOT_ON_DECLS(VectorLength)
+CLAUSE_NOT_ON_DECLS(Wait)
+CLAUSE_NOT_ON_DECLS(Worker)
+#undef CLAUSE_NOT_ON_DECLS
+
+void OpenACCDeclClauseInstantiator::VisitCopyClause(
+ const OpenACCCopyClause &C) {
+ ParsedClause.setVarListDetails(VisitVarList(C.getVarList()),
+ /*IsReadOnly=*/false, /*IsZero=*/false);
+ if (SemaRef.OpenACC().CheckDeclareClause(ParsedClause))
+ return;
+ NewClause = OpenACCCopyClause::Create(
+ SemaRef.getASTContext(), ParsedClause.getClauseKind(),
+ ParsedClause.getBeginLoc(), ParsedClause.getLParenLoc(),
+ ParsedClause.getVarList(), ParsedClause.getEndLoc());
+}
+
+void OpenACCDeclClauseInstantiator::VisitLinkClause(
+ const OpenACCLinkClause &C) {
+ ParsedClause.setVarListDetails(
+ SemaRef.OpenACC().CheckLinkClauseVarList(VisitVarList(C.getVarList())),
+ /*IsReadOnly=*/false, /*IsZero=*/false);
+
+ if (SemaRef.OpenACC().CheckDeclareClause(ParsedClause))
+ return;
+
+ NewClause = OpenACCLinkClause::Create(
+ SemaRef.getASTContext(), ParsedClause.getBeginLoc(),
+ ParsedClause.getLParenLoc(), ParsedClause.getVarList(),
+ ParsedClause.getEndLoc());
+}
+
+void OpenACCDeclClauseInstantiator::VisitDeviceResidentClause(
+ const OpenACCDeviceResidentClause &C) {
+ ParsedClause.setVarListDetails(VisitVarList(C.getVarList()),
+ /*IsReadOnly=*/false, /*IsZero=*/false);
+ if (SemaRef.OpenACC().CheckDeclareClause(ParsedClause))
+ return;
+ NewClause = OpenACCDeviceResidentClause::Create(
+ SemaRef.getASTContext(), ParsedClause.getBeginLoc(),
+ ParsedClause.getLParenLoc(), ParsedClause.getVarList(),
+ ParsedClause.getEndLoc());
+}
+
+void OpenACCDeclClauseInstantiator::VisitCopyInClause(
+ const OpenACCCopyInClause &C) {
+ ParsedClause.setVarListDetails(VisitVarList(C.getVarList()), C.isReadOnly(),
+ /*IsZero=*/false);
+
+ if (SemaRef.OpenACC().CheckDeclareClause(ParsedClause))
+ return;
+ NewClause = OpenACCCopyInClause::Create(
+ SemaRef.getASTContext(), ParsedClause.getClauseKind(),
+ ParsedClause.getBeginLoc(), ParsedClause.getLParenLoc(),
+ ParsedClause.isReadOnly(), ParsedClause.getVarList(),
+ ParsedClause.getEndLoc());
+}
+void OpenACCDeclClauseInstantiator::VisitCopyOutClause(
+ const OpenACCCopyOutClause &C) {
+ ParsedClause.setVarListDetails(VisitVarList(C.getVarList()),
+ /*IsReadOnly=*/false, C.isZero());
+
+ if (SemaRef.OpenACC().CheckDeclareClause(ParsedClause))
+ return;
+ NewClause = OpenACCCopyOutClause::Create(
+ SemaRef.getASTContext(), ParsedClause.getClauseKind(),
+ ParsedClause.getBeginLoc(), ParsedClause.getLParenLoc(),
+ ParsedClause.isZero(), ParsedClause.getVarList(),
+ ParsedClause.getEndLoc());
+}
+void OpenACCDeclClauseInstantiator::VisitCreateClause(
+ const OpenACCCreateClause &C) {
+ ParsedClause.setVarListDetails(VisitVarList(C.getVarList()),
+ /*IsReadOnly=*/false, C.isZero());
+
+ if (SemaRef.OpenACC().CheckDeclareClause(ParsedClause))
+ return;
+ NewClause = OpenACCCreateClause::Create(
+ SemaRef.getASTContext(), ParsedClause.getClauseKind(),
+ ParsedClause.getBeginLoc(), ParsedClause.getLParenLoc(),
+ ParsedClause.isZero(), ParsedClause.getVarList(),
+ ParsedClause.getEndLoc());
+}
+void OpenACCDeclClauseInstantiator::VisitPresentClause(
+ const OpenACCPresentClause &C) {
+ ParsedClause.setVarListDetails(VisitVarList(C.getVarList()),
+ /*IsReadOnly=*/false, /*IsZero=*/false);
+ if (SemaRef.OpenACC().CheckDeclareClause(ParsedClause))
+ return;
+ NewClause = OpenACCPresentClause::Create(
+ SemaRef.getASTContext(), ParsedClause.getBeginLoc(),
+ ParsedClause.getLParenLoc(), ParsedClause.getVarList(),
+ ParsedClause.getEndLoc());
+}
+void OpenACCDeclClauseInstantiator::VisitDevicePtrClause(
+ const OpenACCDevicePtrClause &C) {
+ llvm::SmallVector<Expr *> VarList = VisitVarList(C.getVarList());
+ // Ensure each var is a pointer type.
+ VarList.erase(std::remove_if(VarList.begin(), VarList.end(),
+ [&](Expr *E) {
+ return SemaRef.OpenACC().CheckVarIsPointerType(
+ OpenACCClauseKind::DevicePtr, E);
+ }),
+ VarList.end());
+ ParsedClause.setVarListDetails(VarList,
+ /*IsReadOnly=*/false, /*IsZero=*/false);
+ if (SemaRef.OpenACC().CheckDeclareClause(ParsedClause))
+ return;
+ NewClause = OpenACCDevicePtrClause::Create(
+ SemaRef.getASTContext(), ParsedClause.getBeginLoc(),
+ ParsedClause.getLParenLoc(), ParsedClause.getVarList(),
+ ParsedClause.getEndLoc());
+}
+
+llvm::SmallVector<OpenACCClause *> InstantiateOpenACCClauseList(
+ Sema &S, const MultiLevelTemplateArgumentList &MLTAL,
+ OpenACCDirectiveKind DK, ArrayRef<const OpenACCClause *> ClauseList) {
+ llvm::SmallVector<OpenACCClause *> TransformedClauses;
+
+ for (const auto *Clause : ClauseList) {
+ SemaOpenACC::OpenACCParsedClause ParsedClause(DK, Clause->getClauseKind(),
+ Clause->getBeginLoc());
+ ParsedClause.setEndLoc(Clause->getEndLoc());
+ if (const auto *WithParms = dyn_cast<OpenACCClauseWithParams>(Clause))
+ ParsedClause.setLParenLoc(WithParms->getLParenLoc());
+
+ OpenACCDeclClauseInstantiator Instantiator{S, MLTAL, ParsedClause};
+ Instantiator.Visit(Clause);
+ if (Instantiator.CreatedClause())
+ TransformedClauses.push_back(Instantiator.CreatedClause());
+ }
+ return TransformedClauses;
+}
+
+} // namespace
+
+Decl *TemplateDeclInstantiator::VisitOpenACCDeclareDecl(OpenACCDeclareDecl *D) {
+ SemaRef.OpenACC().ActOnConstruct(D->getDirectiveKind(), D->getBeginLoc());
+ llvm::SmallVector<OpenACCClause *> TransformedClauses =
+ InstantiateOpenACCClauseList(SemaRef, TemplateArgs, D->getDirectiveKind(),
+ D->clauses());
+
+ if (SemaRef.OpenACC().ActOnStartDeclDirective(D->getDirectiveKind(),
+ D->getBeginLoc()))
+ return nullptr;
+
+ DeclGroupRef Res = SemaRef.OpenACC().ActOnEndDeclDirective(
+ D->getDirectiveKind(), D->getBeginLoc(), D->getDirectiveLoc(),
+ D->getEndLoc(), TransformedClauses);
+
+ if (Res.isNull())
+ return nullptr;
+
+ return Res.getSingleDecl();
+}
+
Decl *
TemplateDeclInstantiator::VisitNamespaceAliasDecl(NamespaceAliasDecl *D) {
NamespaceAliasDecl *Inst
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index eaabfae2409f4..ba56179cad7b1 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -11633,7 +11633,8 @@ class OpenACCClauseTransform final
if (!Res.isUsable())
continue;
- Res = Self.getSema().OpenACC().ActOnVar(ParsedClause.getClauseKind(),
+ Res = Self.getSema().OpenACC().ActOnVar(ParsedClause.getDirectiveKind(),
+ ParsedClause.getClauseKind(),
Res.get());
if (Res.isUsable())
@@ -11698,7 +11699,8 @@ void OpenACCClauseTransform<Derived>::VisitSelfClause(
if (!Res.isUsable())
continue;
- Res = Self.getSema().OpenACC().ActOnVar(ParsedClause.getClauseKind(),
+ Res = Self.getSema().OpenACC().ActOnVar(ParsedClause.getDirectiveKind(),
+ ParsedClause.getClauseKind(),
Res.get());
if (Res.isUsable())
@@ -11844,6 +11846,18 @@ void OpenACCClauseTransform<Derived>::VisitCopyClause(
ParsedClause.getVarList(), ParsedClause.getEndLoc());
}
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitLinkClause(
+ const OpenACCLinkClause &C) {
+ llvm_unreachable("link clause not valid unless a decl transform");
+}
+
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitDeviceResidentClause(
+ const OpenACCDeviceResidentClause &C) {
+ llvm_unreachable("device_resident clause not valid unless a decl transform");
+}
+
template <typename Derived>
void OpenACCClauseTransform<Derived>::VisitCopyInClause(
const OpenACCCopyInClause &C) {
diff --git a/clang/lib/Serialization/ASTCommon.cpp b/clang/lib/Serialization/ASTCommon.cpp
index 3a62c4ea5595b..ae8d0a61412b4 100644
--- a/clang/lib/Serialization/ASTCommon.cpp
+++ b/clang/lib/Serialization/ASTCommon.cpp
@@ -458,6 +458,7 @@ bool serialization::isRedeclarableDeclKind(unsigned Kind) {
case Decl::RequiresExprBody:
case Decl::UnresolvedUsingIfExists:
case Decl::HLSLBuffer:
+ case Decl::OpenACCDeclare:
return false;
// These indirectly derive from Redeclarable<T> but are not actually
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index 294e8e063e0a3..ea4b233a6c573 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -12515,7 +12515,7 @@ SmallVector<Expr *> ASTRecordReader::readOpenACCVarList() {
unsigned NumVars = readInt();
llvm::SmallVector<Expr *> VarList;
for (unsigned I = 0; I < NumVars; ++I)
- VarList.push_back(readSubExpr());
+ VarList.push_back(readExpr());
return VarList;
}
@@ -12781,10 +12781,20 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
return OpenACCVectorClause::Create(getContext(), BeginLoc, LParenLoc,
VectorExpr, EndLoc);
}
+ case OpenACCClauseKind::Link: {
+ SourceLocation LParenLoc = readSourceLocation();
+ llvm::SmallVector<Expr *> VarList = readOpenACCVarList();
+ return OpenACCLinkClause::Create(getContext(), BeginLoc, LParenLoc, VarList,
+ EndLoc);
+ }
+ case OpenACCClauseKind::DeviceResident: {
+ SourceLocation LParenLoc = readSourceLocation();
+ llvm::SmallVector<Expr *> VarList = readOpenACCVarList();
+ return OpenACCDeviceResidentClause::Create(getContext(), BeginLoc,
+ LParenLoc, VarList, EndLoc);
+ }
case OpenACCClauseKind::NoHost:
- case OpenACCClauseKind::DeviceResident:
- case OpenACCClauseKind::Link:
case OpenACCClauseKind::Bind:
case OpenACCClauseKind::Invalid:
llvm_unreachable("Clause serialization not yet implemented");
diff --git a/clang/lib/Serialization/ASTReaderDecl.cpp b/clang/lib/Serialization/ASTReaderDecl.cpp
index 17a41fff2267c..2a580c44b94e5 100644
--- a/clang/lib/Serialization/ASTReaderDecl.cpp
+++ b/clang/lib/Serialization/ASTReaderDecl.cpp
@@ -414,6 +414,8 @@ class ASTDeclReader : public DeclVisitor<ASTDeclReader, void> {
void VisitEmptyDecl(EmptyDecl *D);
void VisitLifetimeExtendedTemporaryDecl(LifetimeExtendedTemporaryDecl *D);
+ void VisitOpenACCDeclareDecl(OpenACCDeclareDecl *D);
+
void VisitDeclContext(DeclContext *DC, uint64_t &LexicalOffset,
uint64_t &VisibleOffset, uint64_t &ModuleLocalOffset,
uint64_t &TULocalOffset);
@@ -3099,6 +3101,14 @@ void ASTDeclReader::VisitOMPCapturedExprDecl(OMPCapturedExprDecl *D) {
VisitVarDecl(D);
}
+void ASTDeclReader::VisitOpenACCDeclareDecl(OpenACCDeclareDecl *D) {
+ VisitDecl(D);
+ D->DirKind = Record.readEnum<OpenACCDirectiveKind>();
+ D->DirectiveLoc = Record.readSourceLocation();
+ D->EndLoc = Record.readSourceLocation();
+ Record.readOpenACCClauseList(D->Clauses);
+}
+
//===----------------------------------------------------------------------===//
// Attribute Reading
//===----------------------------------------------------------------------===//
@@ -4204,6 +4214,9 @@ Decl *ASTReader::ReadDeclRecord(GlobalDeclID ID) {
D = ImplicitConceptSpecializationDecl::CreateDeserialized(Context, ID,
Record.readInt());
break;
+ case DECL_OPENACC_DECLARE:
+ D = OpenACCDeclareDecl::CreateDeserialized(Context, ID, Record.readInt());
+ break;
}
assert(D && "Unknown declaration reading AST file");
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index a01e72f193cf1..e51453881a94f 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -1135,6 +1135,7 @@ void ASTWriter::WriteBlockInfoBlock() {
RECORD(DECL_OMP_DECLARE_REDUCTION);
RECORD(DECL_OMP_ALLOCATE);
RECORD(DECL_HLSL_BUFFER);
+ RECORD(DECL_OPENACC_DECLARE);
// Statements and Exprs can occur in the Decls and Types block.
AddStmtsExprs(Stream, Record);
@@ -8828,10 +8829,20 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
AddStmt(const_cast<Expr *>(VC->getIntExpr()));
return;
}
+ case OpenACCClauseKind::Link: {
+ const auto *LC = cast<OpenACCLinkClause>(C);
+ writeSourceLocation(LC->getLParenLoc());
+ writeOpenACCVarList(LC);
+ return;
+ }
+ case OpenACCClauseKind::DeviceResident: {
+ const auto *DRC = cast<OpenACCDeviceResidentClause>(C);
+ writeSourceLocation(DRC->getLParenLoc());
+ writeOpenACCVarList(DRC);
+ return;
+ }
case OpenACCClauseKind::NoHost:
- case OpenACCClauseKind::DeviceResident:
- case OpenACCClauseKind::Link:
case OpenACCClauseKind::Bind:
case OpenACCClauseKind::Invalid:
llvm_unreachable("Clause serialization not yet implemented");
diff --git a/clang/lib/Serialization/ASTWriterDecl.cpp b/clang/lib/Serialization/ASTWriterDecl.cpp
index ac80bb46afa2d..a1810003f5425 100644
--- a/clang/lib/Serialization/ASTWriterDecl.cpp
+++ b/clang/lib/Serialization/ASTWriterDecl.cpp
@@ -176,6 +176,8 @@ namespace clang {
void VisitOMPDeclareMapperDecl(OMPDeclareMapperDecl *D);
void VisitOMPCapturedExprDecl(OMPCapturedExprDecl *D);
+ void VisitOpenACCDeclareDecl(OpenACCDeclareDecl *D);
+
/// Add an Objective-C type parameter list to the given record.
void AddObjCTypeParamList(ObjCTypeParamList *typeParams) {
// Empty type parameter list.
@@ -2258,6 +2260,16 @@ void ASTDeclWriter::VisitOMPCapturedExprDecl(OMPCapturedExprDecl *D) {
Code = serialization::DECL_OMP_CAPTUREDEXPR;
}
+void ASTDeclWriter::VisitOpenACCDeclareDecl(OpenACCDeclareDecl *D) {
+ Record.writeUInt32(D->clauses().size());
+ VisitDecl(D);
+ Record.writeEnum(D->DirKind);
+ Record.AddSourceLocation(D->DirectiveLoc);
+ Record.AddSourceLocation(D->EndLoc);
+ Record.writeOpenACCClauseList(D->clauses());
+ Code = serialization::DECL_OPENACC_DECLARE;
+}
+
//===----------------------------------------------------------------------===//
// ASTWriter Implementation
//===----------------------------------------------------------------------===//
diff --git a/clang/test/AST/ast-print-openacc-declare-construct.cpp b/clang/test/AST/ast-print-openacc-declare-construct.cpp
new file mode 100644
index 0000000000000..fce4afc6aedae
--- /dev/null
+++ b/clang/test/AST/ast-print-openacc-declare-construct.cpp
@@ -0,0 +1,52 @@
+// RUN: %clang_cc1 -fopenacc -ast-print %s -o - | FileCheck %s
+
+int *Global, *Global2;
+int GlobalArray[5];
+int GlobalArray2[5];
+// CHECK: #pragma acc declare deviceptr(Global) copyin(GlobalArray)
+#pragma acc declare deviceptr(Global), copyin(GlobalArray)
+// CHECK: #pragma acc declare create(Global2, GlobalArray2)
+#pragma acc declare create(Global2, GlobalArray2)
+
+namespace NS {
+int NSVar;
+int NSArray[5];
+// CHECK: #pragma acc declare create(NSVar, NSArray)
+#pragma acc declare create(NSVar, NSArray)
+}
+
+struct Struct {
+ static const int StaticMem = 5;
+ static const int StaticMemArray[5];
+// CHECK: #pragma acc declare copyin(StaticMem, StaticMemArray)
+#pragma acc declare copyin(StaticMem, StaticMemArray)
+
+ void MemFunc1(int Arg) {
+ int Local;
+ int LocalArray[5];
+// CHECK: #pragma acc declare present(Arg, Local, LocalArray)
+#pragma acc declare present(Arg, Local, LocalArray)
+ }
+ void MemFunc2(int Arg);
+};
+void Struct::MemFunc2(int Arg) {
+ int Local;
+ int LocalArray[5];
+// CHECK: #pragma acc declare present(Arg, Local, LocalArray)
+#pragma acc declare present(Arg, Local, LocalArray)
+}
+
+void NormalFunc(int Arg) {
+ int Local;
+ int LocalArray[5];
+// CHECK: #pragma acc declare present(Arg, Local, LocalArray)
+#pragma acc declare present(Arg, Local, LocalArray)
+}
+
+void NormalFunc2(int *Arg) {
+ int Local;
+ int LocalArray[5];
+ extern int ExternLocal;
+// CHECK: #pragma acc declare deviceptr(Arg) device_resident(Local) link(ExternLocal)
+#pragma acc declare deviceptr(Arg) device_resident(Local) link(ExternLocal)
+}
diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c
index b871624b6e943..7c987ba0a1bd7 100644
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@@ -520,23 +520,25 @@ void VarListClauses() {
#pragma acc exit data delete(s.array[s.value : 5], s.value),async
for(int i = 0; i < 5;++i) {}
- // expected-error at +2{{expected ','}}
- // expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented, clause ignored}}
-#pragma acc serial device_resident(s.array[s.value] s.array[s.value :5] ), self
- for(int i = 0; i < 5;++i) {}
+ // expected-error at +3{{expected ','}}
+ // expected-error at +2{{OpenACC variable on 'declare' construct is not a valid variable name or array name}}
+ // expected-error at +1{{OpenACC variable on 'declare' construct is not a valid variable name or array name}}
+#pragma acc declare device_resident(s.array[s.value] s.array[s.value :5] ), copy(s)
- // expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented, clause ignored}}
-#pragma acc serial device_resident(s.array[s.value : 5], s.value), self
- for(int i = 0; i < 5;++i) {}
+ int CopyRef1, CopyRef2, CopyRef3;
- // expected-error at +2{{expected ','}}
- // expected-warning at +1{{OpenACC clause 'link' not yet implemented, clause ignored}}
-#pragma acc serial link(s.array[s.value] s.array[s.value :5] ), self
- for(int i = 0; i < 5;++i) {}
+ // expected-error at +2{{OpenACC variable on 'declare' construct is not a valid variable name or array name}}
+ // expected-error at +1{{OpenACC variable on 'declare' construct is not a valid variable name or array name}}
+#pragma acc declare device_resident(s.array[s.value : 5], s.value), copy(CopyRef1)
- // expected-warning at +1{{OpenACC clause 'link' not yet implemented, clause ignored}}
-#pragma acc serial link(s.array[s.value : 5], s.value), self
- for(int i = 0; i < 5;++i) {}
+ // expected-error at +3{{expected ','}}
+ // expected-error at +2{{OpenACC variable on 'declare' construct is not a valid variable name or array name}}
+ // expected-error at +1{{OpenACC variable on 'declare' construct is not a valid variable name or array name}}
+#pragma acc declare link(s.array[s.value] s.array[s.value :5] ), copy(CopyRef2)
+
+ // expected-error at +2{{OpenACC variable on 'declare' construct is not a valid variable name or array name}}
+ // expected-error at +1{{OpenACC variable on 'declare' construct is not a valid variable name or array name}}
+#pragma acc declare link(s.array[s.value : 5], s.value), copy(CopyRef3)
// expected-error at +1{{expected ','}}
#pragma acc update host(s.array[s.value] s.array[s.value :5] )
diff --git a/clang/test/ParserOpenACC/parse-constructs.c b/clang/test/ParserOpenACC/parse-constructs.c
index f0698495a3cc2..976220979f0b9 100644
--- a/clang/test/ParserOpenACC/parse-constructs.c
+++ b/clang/test/ParserOpenACC/parse-constructs.c
@@ -131,7 +131,7 @@ void func() {
// expected-error at +2{{invalid OpenACC clause 'clause'}}
- // expected-warning at +1{{OpenACC construct 'declare' not yet implemented, pragma ignored}}
+ // expected-error at +1{{no valid clauses specified in OpenACC 'declare' directive}}
#pragma acc declare clause list
for(;;){}
// expected-error at +1{{invalid OpenACC clause 'clause'}}
diff --git a/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c b/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
index 9e74ce27ffbd9..bbc40f4b401b2 100644
--- a/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
+++ b/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
@@ -80,7 +80,7 @@ void uses() {
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop auto deviceptr(VarPtr)
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
+ // expected-error at +1{{OpenACC 'device_resident' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop auto device_resident(VarPtr)
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop auto firstprivate(Var)
@@ -88,7 +88,7 @@ void uses() {
// expected-error at +1{{OpenACC 'host' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop auto host(Var)
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'link' not yet implemented}}
+ // expected-error at +1{{OpenACC 'link' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop auto link(Var)
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop auto no_create(Var)
@@ -197,7 +197,7 @@ void uses() {
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop deviceptr(VarPtr) auto
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
+ // expected-error at +1{{OpenACC 'device_resident' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop device_resident(VarPtr) auto
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop firstprivate(Var) auto
@@ -205,7 +205,7 @@ void uses() {
// expected-error at +1{{OpenACC 'host' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop host(Var) auto
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'link' not yet implemented}}
+ // expected-error at +1{{OpenACC 'link' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop link(Var) auto
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop no_create(Var) auto
@@ -315,7 +315,7 @@ void uses() {
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop independent deviceptr(VarPtr)
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
+ // expected-error at +1{{OpenACC 'device_resident' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop independent device_resident(VarPtr)
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop independent firstprivate(Var)
@@ -323,7 +323,7 @@ void uses() {
// expected-error at +1{{OpenACC 'host' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop independent host(Var)
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'link' not yet implemented}}
+ // expected-error at +1{{OpenACC 'link' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop independent link(Var)
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop independent no_create(Var)
@@ -432,7 +432,7 @@ void uses() {
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop deviceptr(VarPtr) independent
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
+ // expected-error at +1{{OpenACC 'device_resident' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop device_resident(VarPtr) independent
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop firstprivate(Var) independent
@@ -440,7 +440,7 @@ void uses() {
// expected-error at +1{{OpenACC 'host' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop host(Var) independent
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'link' not yet implemented}}
+ // expected-error at +1{{OpenACC 'link' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop link(Var) independent
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop no_create(Var) independent
@@ -558,7 +558,7 @@ void uses() {
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop seq deviceptr(VarPtr)
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
+ // expected-error at +1{{OpenACC 'device_resident' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop seq device_resident(VarPtr)
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop seq firstprivate(Var)
@@ -566,7 +566,7 @@ void uses() {
// expected-error at +1{{OpenACC 'host' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop seq host(Var)
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'link' not yet implemented}}
+ // expected-error at +1{{OpenACC 'link' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop seq link(Var)
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop seq no_create(Var)
@@ -681,7 +681,7 @@ void uses() {
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop deviceptr(VarPtr) seq
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
+ // expected-error at +1{{OpenACC 'device_resident' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop device_resident(VarPtr) seq
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop firstprivate(Var) seq
@@ -689,7 +689,7 @@ void uses() {
// expected-error at +1{{OpenACC 'host' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop host(Var) seq
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'link' not yet implemented}}
+ // expected-error at +1{{OpenACC 'link' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop link(Var) seq
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop no_create(Var) seq
diff --git a/clang/test/SemaOpenACC/combined-construct-device_type-clause.c b/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
index c185f607548ff..2d7a9418cce15 100644
--- a/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
@@ -107,8 +107,7 @@ void uses() {
// expected-note at +1{{previous clause is here}}
#pragma acc serial loop device_type(*) deviceptr(VarPtr)
for(int i = 0; i < 5; ++i);
- // expected-error at +2{{OpenACC clause 'device_resident' may not follow a 'device_type' clause in a 'kernels loop' construct}}
- // expected-note at +1{{previous clause is here}}
+ // expected-error at +1{{OpenACC 'device_resident' clause is not valid on 'kernels loop' directive}}
#pragma acc kernels loop device_type(*) device_resident(VarPtr)
for(int i = 0; i < 5; ++i);
// expected-error at +2{{OpenACC clause 'firstprivate' may not follow a 'device_type' clause in a 'parallel loop' construct}}
@@ -118,8 +117,7 @@ void uses() {
// expected-error at +1{{OpenACC 'host' clause is not valid on 'serial loop' directive}}
#pragma acc serial loop device_type(*) host(Var)
for(int i = 0; i < 5; ++i);
- // expected-error at +2{{OpenACC clause 'link' may not follow a 'device_type' clause in a 'parallel loop' construct}}
- // expected-note at +1{{previous clause is here}}
+ // expected-error at +1{{OpenACC 'link' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop device_type(*) link(Var)
for(int i = 0; i < 5; ++i);
// expected-error at +2{{OpenACC clause 'no_create' may not follow a 'device_type' clause in a 'serial loop' construct}}
diff --git a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
index 4290fb7665685..5f464cc94b092 100644
--- a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
@@ -113,8 +113,7 @@ void uses() {
// expected-note at +1{{previous clause is here}}
#pragma acc kernels device_type(*) deviceptr(VarPtr)
while(1);
- // expected-error at +2{{OpenACC clause 'device_resident' may not follow a 'device_type' clause in a 'kernels' construct}}
- // expected-note at +1{{previous clause is here}}
+ // expected-error at +1{{OpenACC 'device_resident' clause is not valid on 'kernels' directive}}
#pragma acc kernels device_type(*) device_resident(VarPtr)
while(1);
// expected-error at +2{{OpenACC clause 'firstprivate' may not follow a 'device_type' clause in a 'parallel' construct}}
@@ -124,8 +123,7 @@ void uses() {
// expected-error at +1{{OpenACC 'host' clause is not valid on 'kernels' directive}}
#pragma acc kernels device_type(*) host(Var)
while(1);
- // expected-error at +2{{OpenACC clause 'link' may not follow a 'device_type' clause in a 'kernels' construct}}
- // expected-note at +1{{previous clause is here}}
+ // expected-error at +1{{OpenACC 'link' clause is not valid on 'kernels' directive}}
#pragma acc kernels device_type(*) link(Var)
while(1);
// expected-error at +2{{OpenACC clause 'no_create' may not follow a 'device_type' clause in a 'kernels' construct}}
diff --git a/clang/test/SemaOpenACC/data-construct.cpp b/clang/test/SemaOpenACC/data-construct.cpp
index 4c868b68e332e..e1d0c04d7fdee 100644
--- a/clang/test/SemaOpenACC/data-construct.cpp
+++ b/clang/test/SemaOpenACC/data-construct.cpp
@@ -165,7 +165,6 @@ struct HasMembers {
// expected-error at +1{{OpenACC variable in 'use_device' clause is not a valid variable name or array name}}
#pragma acc host_data use_device(this->Member)
;
- // expected-error at +1{{OpenACC variable in 'use_device' clause is not a valid variable name or array name}}
#pragma acc host_data use_device(Member)
;
}
diff --git a/clang/test/SemaOpenACC/declare-construct-ast.cpp b/clang/test/SemaOpenACC/declare-construct-ast.cpp
new file mode 100644
index 0000000000000..dfc79039fc9e6
--- /dev/null
+++ b/clang/test/SemaOpenACC/declare-construct-ast.cpp
@@ -0,0 +1,390 @@
+// 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
+
+int *Global;
+// CHECK: VarDecl{{.*}}Global 'int *'
+int GlobalArray[5];
+// CHECK-NEXT: VarDecl{{.*}}GlobalArray 'int[5]'
+#pragma acc declare deviceptr(Global), copyin(GlobalArray)
+// CHECK-NEXT: OpenACCDeclareDecl
+// CHECK-NEXT: deviceptr clause
+// CHECK-NEXT: DeclRefExpr{{.*}}'Global' 'int *'
+// CHECK-NEXT: copyin clause
+// CHECK-NEXT: DeclRefExpr{{.*}}'GlobalArray' 'int[5]'
+
+int *Global2;
+// CHECK: VarDecl{{.*}}Global2 'int *'
+int GlobalArray2[5];
+// CHECK-NEXT: VarDecl{{.*}}GlobalArray2 'int[5]'
+#pragma acc declare create(Global2, GlobalArray2)
+// CHECK-NEXT: OpenACCDeclareDecl
+// CHECK-NEXT: create clause
+// CHECK-NEXT: DeclRefExpr{{.*}}'Global2' 'int *'
+// CHECK-NEXT: DeclRefExpr{{.*}}'GlobalArray2' 'int[5]'
+
+int Global3;
+// CHECK: VarDecl{{.*}}Global3 'int'
+int GlobalArray3[5];
+// CHECK-NEXT: VarDecl{{.*}}GlobalArray3 'int[5]'
+#pragma acc declare link(Global3) device_resident(GlobalArray3)
+// CHECK-NEXT: OpenACCDeclareDecl
+// CHECK-NEXT: link clause
+// CHECK-NEXT: DeclRefExpr{{.*}}'Global3' 'int'
+// CHECK-NEXT: device_resident clause
+// CHECK-NEXT: DeclRefExpr{{.*}}'GlobalArray3' 'int[5]'
+
+namespace NS {
+int NSVar;
+// CHECK: VarDecl{{.*}}NSVar 'int'
+int NSArray[5];
+// CHECK-NEXT: VarDecl{{.*}}NSArray 'int[5]'
+#pragma acc declare create(NSVar, NSArray)
+// CHECK-NEXT: OpenACCDeclareDecl
+// CHECK-NEXT: create clause
+// CHECK-NEXT: DeclRefExpr{{.*}}'NSVar' 'int'
+// CHECK-NEXT: DeclRefExpr{{.*}}'NSArray' 'int[5]'
+}
+
+struct Struct {
+ // CHECK-NEXT: CXXRecordDecl{{.*}} Struct definition
+ // Skip DefinitionData and go right to the definition.
+ // CHECK: CXXRecordDecl{{.*}} implicit struct Struct
+ static const int StaticMem = 5;
+ // CHECK-NEXT: VarDecl {{.*}} StaticMem 'const int' static cinit
+ // CHECK-NEXT: IntegerLiteral{{.*}}'int' 5
+ static const int StaticMem2 = 5;
+ // CHECK-NEXT: VarDecl {{.*}} StaticMem2 'const int' static cinit
+ // CHECK-NEXT: IntegerLiteral{{.*}}'int' 5
+ static const int StaticMemArray[5];
+ // CHECK-NEXT: VarDecl {{.*}} StaticMemArray 'const int[5]' static
+ static const int StaticMemArray2[5];
+ // CHECK-NEXT: VarDecl {{.*}} StaticMemArray2 'const int[5]' static
+#pragma acc declare copyin(StaticMem, StaticMemArray) create(StaticMem2, StaticMemArray2)
+ // CHECK-NEXT: OpenACCDeclareDecl
+ // CHECK-NEXT: copyin clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'StaticMem' 'const int'
+ // CHECK-NEXT: DeclRefExpr{{.*}}'StaticMemArray' 'const int[5]'
+ // CHECK-NEXT: create clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'StaticMem2' 'const int'
+ // CHECK-NEXT: DeclRefExpr{{.*}}'StaticMemArray2' 'const int[5]'
+
+ void MemFunc1(int Arg) {
+ // CHECK-NEXT: CXXMethodDecl{{.*}}MemFunc1 'void (int)'
+ // CHECK-NEXT: ParmVarDecl{{.*}} Arg 'int'
+ // CHECK-NEXT: CompoundStmt
+ int Local;
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} Local 'int'
+ int LocalArray[5];
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} LocalArray 'int[5]'
+#pragma acc declare present(Arg, Local, LocalArray)
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: OpenACCDeclareDecl
+ // CHECK-NEXT: present clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Arg' 'int'
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Local' 'int'
+ // CHECK-NEXT: DeclRefExpr{{.*}}'LocalArray' 'int[5]'
+ }
+ void MemFunc2(int Arg);
+ // CHECK: CXXMethodDecl{{.*}}MemFunc2
+};
+void Struct::MemFunc2(int Arg) {
+ // CHECK: CXXMethodDecl{{.*}}MemFunc2 'void (int)'
+ // CHECK-NEXT: ParmVarDecl{{.*}} Arg 'int'
+ // CHECK-NEXT: CompoundStmt
+ int Local;
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} Local 'int'
+ int LocalArray[5];
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} LocalArray 'int[5]'
+#pragma acc declare present(Arg, Local, LocalArray)
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: OpenACCDeclareDecl
+ // CHECK-NEXT: present clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Arg' 'int'
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Local' 'int'
+ // CHECK-NEXT: DeclRefExpr{{.*}}'LocalArray' 'int[5]'
+}
+
+void NormalFunc(int Arg) {
+ // CHECK-NEXT: FunctionDecl{{.*}}NormalFunc 'void (int)'
+ // CHECK-NEXT: ParmVarDecl{{.*}} Arg 'int'
+ // CHECK-NEXT: CompoundStmt
+ int Local;
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} Local 'int'
+ int LocalArray[5];
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} LocalArray 'int[5]'
+#pragma acc declare present(Arg, Local, LocalArray)
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: OpenACCDeclareDecl
+ // CHECK-NEXT: present clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Arg' 'int'
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Local' 'int'
+ // CHECK-NEXT: DeclRefExpr{{.*}}'LocalArray' 'int[5]'
+}
+
+template<typename T>
+struct DependentStruct {
+ // CHECK: ClassTemplateDecl{{.*}}DependentStruct
+ // CHECK-NEXT: TemplateTypeParmDecl{{.*}}depth 0 index 0 T
+ // CHECK-NEXT: CXXRecordDecl{{.*}}DependentStruct definition
+ // CHECK: CXXRecordDecl{{.*}}implicit struct DependentStruct
+ static const T StaticMem = 5;
+ // CHECK-NEXT: VarDecl{{.*}} StaticMem 'const T' static cinit
+ // CHECK-NEXT: IntegerLiteral{{.*}}'int' 5
+ static const T StaticMem2 = 5;
+ // CHECK-NEXT: VarDecl{{.*}} StaticMem2 'const T' static cinit
+ // CHECK-NEXT: IntegerLiteral{{.*}}'int' 5
+ static constexpr T StaticMemArray[5] = {};
+ // CHECK-NEXT: VarDecl{{.*}} StaticMemArray 'const T[5]'
+ // CHECK-NEXT: InitListExpr{{.*}}'void'
+ static constexpr T StaticMemArray2[5] = {};
+ // CHECK-NEXT: VarDecl{{.*}} StaticMemArray2 'const T[5]'
+ // CHECK-NEXT: InitListExpr{{.*}}'void'
+#pragma acc declare copyin(StaticMem, StaticMemArray) create(StaticMem2, StaticMemArray2)
+ // CHECK-NEXT: OpenACCDeclareDecl
+ // CHECK-NEXT: copyin clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'StaticMem' 'const T'
+ // CHECK-NEXT: DeclRefExpr{{.*}}'StaticMemArray' 'const T[5]'
+ // CHECK-NEXT: create clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'StaticMem2' 'const T'
+ // CHECK-NEXT: DeclRefExpr{{.*}}'StaticMemArray2' 'const T[5]'
+
+ template<typename U>
+ void DepMemFunc1(U Arg, U Arg2) {
+ // CHECK-NEXT: FunctionTemplateDecl{{.*}}DepMemFunc1
+ // CHECK-NEXT: TemplateTypeParmDecl{{.*}}depth 1 index 0 U
+ // CHECK-NEXT: CXXMethodDecl{{.*}}DepMemFunc1 'void (U, U)'
+ // CHECK-NEXT: ParmVarDecl{{.*}} Arg 'U'
+ // CHECK-NEXT: ParmVarDecl{{.*}} Arg2 'U'
+ // CHECK-NEXT: CompoundStmt
+ T Local, Local2;
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} Local 'T'
+ // CHECK-NEXT: VarDecl{{.*}} Local2 'T'
+ U LocalArray[5];
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} LocalArray 'U[5]'
+ U LocalArray2[5];
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} LocalArray2 'U[5]'
+#pragma acc declare copy(Arg, Local, LocalArray) copyout(Arg2, Local2, LocalArray2)
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: OpenACCDeclareDecl
+ // CHECK-NEXT: copy clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Arg' 'U'
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Local' 'T'
+ // CHECK-NEXT: DeclRefExpr{{.*}}'LocalArray' 'U[5]'
+ // CHECK-NEXT: copyout clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Arg2' 'U'
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Local2' 'T'
+ // CHECK-NEXT: DeclRefExpr{{.*}}'LocalArray2' 'U[5]'
+
+ extern T Local3;
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} Local3 'T' extern
+ T Local4;
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} Local4 'T'
+#pragma acc declare link(Local3) device_resident(Local4)
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: OpenACCDeclareDecl
+ // CHECK-NEXT: link clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Local3' 'T'
+ // CHECK-NEXT: device_resident clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Local4' 'T'
+ }
+ template<typename U>
+ void DepMemFunc2(U Arg);
+ // CHECK-NEXT: FunctionTemplateDecl{{.*}}DepMemFunc2
+ // CHECK-NEXT: TemplateTypeParmDecl{{.*}}depth 1 index 0 U
+ // CHECK-NEXT: CXXMethodDecl{{.*}}DepMemFunc2 'void (U)'
+ // CHECK-NEXT: ParmVarDecl{{.*}} Arg 'U'
+};
+
+// Instantiation of class.
+// CHECK-NEXT: ClassTemplateSpecializationDecl{{.*}}DependentStruct definition
+// CHECK: TemplateArgument type 'int'
+// CHECK-NEXT: BuiltinType{{.*}}'int'
+// CHECK-NEXT: CXXRecordDecl{{.*}} struct DependentStruct
+
+// CHECK-NEXT: VarDecl{{.*}} StaticMem 'const int'
+// CHECK-NEXT: IntegerLiteral{{.*}}'int' 5
+
+// CHECK-NEXT: VarDecl{{.*}} StaticMem2 'const int'
+// CHECK-NEXT: IntegerLiteral{{.*}}'int' 5
+//
+// CHECK-NEXT: VarDecl{{.*}} StaticMemArray 'const int[5]'
+// CHECK-NEXT: value: Array size=5
+// CHECK-NEXT: filler: 5 x Int 0
+// CHECK-NEXT: InitListExpr{{.*}} 'const int[5]'
+// CHECK-NEXT: array_filler
+
+// CHECK-NEXT: VarDecl{{.*}} StaticMemArray2 'const int[5]'
+// CHECK-NEXT: value: Array size=5
+// CHECK-NEXT: filler: 5 x Int 0
+// CHECK-NEXT: InitListExpr{{.*}} 'const int[5]'
+// CHECK-NEXT: array_filler
+
+// CHECK-NEXT: OpenACCDeclareDecl
+// CHECK-NEXT: copyin clause
+// CHECK-NEXT: DeclRefExpr{{.*}}'StaticMem' 'const int'
+// CHECK-NEXT: DeclRefExpr{{.*}}'StaticMemArray' 'const int[5]'
+// CHECK-NEXT: create clause
+// CHECK-NEXT: DeclRefExpr{{.*}}'StaticMem2' 'const int'
+// CHECK-NEXT: DeclRefExpr{{.*}}'StaticMemArray2' 'const int[5]'
+
+// CHECK-NEXT: FunctionTemplateDecl{{.*}} DepMemFunc1
+// CHECK-NEXT: TemplateTypeParmDecl{{.*}}depth 0 index 0 U
+// CHECK-NEXT: CXXMethodDecl{{.*}}DepMemFunc1 'void (U, U)'
+// CHECK-NEXT: ParmVarDecl{{.*}} Arg 'U'
+// CHECK-NEXT: ParmVarDecl{{.*}} Arg2 'U'
+// CHECK-NEXT: CXXMethodDecl{{.*}}DepMemFunc1 'void (float, float)'
+// CHECK-NEXT: TemplateArgument type 'float'
+// CHECK-NEXT: BuiltinType{{.*}}'float'
+// CHECK-NEXT: ParmVarDecl{{.*}} Arg 'float'
+// CHECK-NEXT: ParmVarDecl{{.*}} Arg2 'float'
+// CHECK-NEXT: CompoundStmt
+
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT: VarDecl{{.*}} Local 'int'
+// CHECK-NEXT: VarDecl{{.*}} Local2 'int'
+
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT: VarDecl{{.*}} LocalArray 'float[5]'
+
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT: VarDecl{{.*}} LocalArray2 'float[5]'
+
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT: OpenACCDeclareDecl
+// CHECK-NEXT: copy clause
+// CHECK-NEXT: DeclRefExpr{{.*}}'Arg' 'float'
+// CHECK-NEXT: DeclRefExpr{{.*}}'Local' 'int'
+// CHECK-NEXT: DeclRefExpr{{.*}}'LocalArray' 'float[5]'
+// CHECK-NEXT: copyout clause
+// CHECK-NEXT: DeclRefExpr{{.*}}'Arg2' 'float'
+// CHECK-NEXT: DeclRefExpr{{.*}}'Local2' 'int'
+// CHECK-NEXT: DeclRefExpr{{.*}}'LocalArray2' 'float[5]'
+
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT: VarDecl{{.*}} Local3 'int' extern
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT: VarDecl{{.*}} Local4 'int'
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT: OpenACCDeclareDecl
+// CHECK-NEXT: link clause
+// CHECK-NEXT: DeclRefExpr{{.*}}'Local3' 'int'
+// CHECK-NEXT: device_resident clause
+// CHECK-NEXT: DeclRefExpr{{.*}}'Local4' 'int'
+
+// CHECK-NEXT: FunctionTemplateDecl{{.*}}DepMemFunc2
+// CHECK-NEXT: TemplateTypeParmDecl{{.*}}depth 0 index 0 U
+// CHECK-NEXT: CXXMethodDecl{{.*}}DepMemFunc2 'void (U)'
+// CHECK-NEXT: ParmVarDecl{{.*}} Arg 'U'
+// CHECK-NEXT: CXXMethodDecl{{.*}}DepMemFunc2 'void (float)'
+// CHECK-NEXT: TemplateArgument type 'float'
+// CHECK-NEXT: BuiltinType{{.*}}'float'
+// CHECK-NEXT: ParmVarDecl{{.*}} Arg 'float'
+// CHECK-NEXT: CompoundStmt
+
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT: VarDecl{{.*}} Local 'int'
+
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT: VarDecl{{.*}} LocalArray 'float[5]'
+
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT: OpenACCDeclareDecl
+// CHECK-NEXT: present clause
+// CHECK-NEXT: DeclRefExpr{{.*}}'Arg' 'float'
+// CHECK-NEXT: DeclRefExpr{{.*}}'Local' 'int'
+// CHECK-NEXT: DeclRefExpr{{.*}}'LocalArray' 'float[5]'
+
+template<typename T>
+template<typename U>
+void DependentStruct<T>::DepMemFunc2(U Arg) {
+ // CHECK: FunctionTemplateDecl{{.*}} DepMemFunc2
+ // CHECK-NEXT: TemplateTypeParmDecl{{.*}}depth 1 index 0 U
+ // CHECK-NEXT: CXXMethodDecl{{.*}}DepMemFunc2 'void (U)'
+ // CHECK-NEXT: ParmVarDecl{{.*}} Arg 'U'
+ // CHECK-NEXT: CompoundStmt
+ T Local;
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} Local 'T'
+ U LocalArray[5];
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} LocalArray 'U[5]'
+#pragma acc declare present(Arg, Local, LocalArray)
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: OpenACCDeclareDecl
+ // CHECK-NEXT: present clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Arg' 'U'
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Local' 'T'
+ // CHECK-NEXT: DeclRefExpr{{.*}}'LocalArray' 'U[5]'
+}
+
+template<typename T, unsigned Size>
+void DependentFunc(T Arg) {
+ // CHECK: FunctionTemplateDecl{{.*}} DependentFunc
+ // CHECK-NEXT: TemplateTypeParmDecl{{.*}}depth 0 index 0 T
+ // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} depth 0 index 1 Size
+ // CHECK-NEXT: FunctionDecl{{.*}}DependentFunc 'void (T)'
+ // CHECK-NEXT: ParmVarDecl{{.*}} Arg 'T'
+ // CHECK-NEXT: CompoundStmt
+ T Local;
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} Local 'T'
+ T LocalArray[Size];
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} LocalArray 'T[Size]'
+
+#pragma acc declare present(Arg, Local, LocalArray)
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: OpenACCDeclareDecl
+ // CHECK-NEXT: present clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Arg' 'T'
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Local' 'T'
+ // CHECK-NEXT: DeclRefExpr{{.*}}'LocalArray' 'T[Size]'
+
+ // Instantiation:
+ // CHECK-NEXT: FunctionDecl{{.*}} DependentFunc 'void (int)'
+ // CHECK-NEXT: TemplateArgument type 'int'
+ // CHECK-NEXT: BuiltinType{{.*}}'int'
+ // CHECK-NEXT: TemplateArgument integral '5U'
+ // CHECK-NEXT: ParmVarDecl{{.*}} Arg 'int'
+ // CHECK-NEXT: CompoundStmt
+
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} Local 'int'
+
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} LocalArray 'int[5]'
+
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: OpenACCDeclareDecl
+ // CHECK-NEXT: present clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Arg' 'int'
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Local' 'int'
+ // CHECK-NEXT: DeclRefExpr{{.*}}'LocalArray' 'int[5]'
+}
+
+void use() {
+ float i;
+ DependentStruct<int> S;
+ S.DepMemFunc1(i, i);
+ S.DepMemFunc2(i);
+ DependentFunc<int, 5>(i);
+}
+
+#endif // PCH_HELPER
diff --git a/clang/test/SemaOpenACC/declare-construct.cpp b/clang/test/SemaOpenACC/declare-construct.cpp
new file mode 100644
index 0000000000000..a1fed096635fa
--- /dev/null
+++ b/clang/test/SemaOpenACC/declare-construct.cpp
@@ -0,0 +1,308 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+int *Global;
+int GlobalArray[5];
+// expected-error at +1{{no valid clauses specified in OpenACC 'declare' directive}}
+#pragma acc declare
+namespace NS {
+int *NSVar;
+int NSArray[5];
+// expected-error at +2{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+// expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(Global, GlobalArray)
+// Ok, correct scope.
+#pragma acc declare create(NSVar, NSArray)
+
+// expected-error at +4{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -3{{previous reference is here}}
+// expected-error at +2{{variable referenced in 'copyin' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -5{{previous reference is here}}
+#pragma acc declare create(NSVar) copyin(NSVar)
+
+// expected-error at +1{{no valid clauses specified in OpenACC 'declare' directive}}
+#pragma acc declare
+
+int NSVar1, NSVar2, NSVar3, NSVar4, NSVar5, *NSVar6, NSVar7, NSVar8;
+
+// Only create, copyin, deviceptr, device-resident, link at NS scope.
+// expected-error at +3{{OpenACC 'copy' clause on a 'declare' directive is not allowed at global or namespace scope}}
+// expected-error at +2{{OpenACC 'copyout' clause on a 'declare' directive is not allowed at global or namespace scope}}
+// expected-error at +1{{OpenACC 'present' clause on a 'declare' directive is not allowed at global or namespace scope}}
+#pragma acc declare copy(NSVar1) copyin(NSVar2), copyout(NSVar3), create(NSVar4), present(NSVar5), deviceptr(NSVar6), device_resident(NSVar7), link(NSVar8)
+
+extern "C" {
+ int ExternVar, ExternVar1, ExternVar2, ExternVar3, ExternVar4, *ExternVar5, ExternVar6, ExternVar7;
+ // Only create, copyin, deviceptr, device-resident, link at NS scope.
+ // expected-error at +3{{OpenACC 'copy' clause on a 'declare' directive is not allowed at global or namespace scope}}
+ // expected-error at +2{{OpenACC 'copyout' clause on a 'declare' directive is not allowed at global or namespace scope}}
+ // expected-error at +1{{OpenACC 'present' clause on a 'declare' directive is not allowed at global or namespace scope}}
+#pragma acc declare copy(ExternVar) copyin(ExternVar1), copyout(ExternVar2), create(ExternVar3), present(ExternVar4), deviceptr(ExternVar5), device_resident(ExternVar6), link(ExternVar7)
+ }
+}
+// expected-error at +2{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+// expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(NS::NSVar, NS::NSArray)
+
+struct Struct {
+ static const int StaticMem = 5;
+ static const int StaticMem2 = 5;
+ int NonStaticMem;
+// expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(Global)
+ // OK, same scope.
+#pragma acc declare create(StaticMem, StaticMem2)
+// expected-error at +4{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -2{{previous reference is here}}
+// expected-error at +2{{variable referenced in 'copyin' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -4{{previous reference is here}}
+#pragma acc declare create(StaticMem) copyin(StaticMem)
+// expected-error at +1{{no valid clauses specified in OpenACC 'declare' directive}}
+#pragma acc declare
+
+ void Inline(int Arg) {
+// expected-error at +1{{no valid clauses specified in OpenACC 'declare' directive}}
+#pragma acc declare
+// expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(StaticMem)
+
+ int Local, Local2, Local3, Local4;
+ // OK, same scope.
+#pragma acc declare create(Local, Arg)
+// expected-error at +2{{variable referenced in 'copyin' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at +1{{previous reference is here}}
+#pragma acc declare create(Local2) copyin(Local2)
+
+ for (int I = 0; I < 5; ++I) {
+ int Other;
+ // FIXME: We don't catch this because we use decl-context instead of scope.
+#pragma acc declare create(Local3, Local4)
+ // OK, same scope.
+#pragma acc declare create(I, Other)
+// expected-error at +4 2{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -2 2{{previous reference is here}}
+// expected-error at +2 2{{variable referenced in 'copyin' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -4 2{{previous reference is here}}
+#pragma acc declare create(I, Other) copyin(I, Other)
+
+// expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(NonStaticMem)
+ }
+ }
+
+ void OutOfLine(int Arg, int Arg2);
+};
+
+void Struct::OutOfLine(int Arg, int Arg2) {
+// expected-error at +1{{no valid clauses specified in OpenACC 'declare' directive}}
+#pragma acc declare
+// expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(StaticMem)
+
+ int Local, Local2;
+// OK, same scope.
+#pragma acc declare create(Local, Arg)
+// expected-error at +4{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -2{{previous reference is here}}
+// expected-error at +2{{variable referenced in 'copyin' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -4{{previous reference is here}}
+#pragma acc declare create(Local) copyin(Local)
+
+ for (int I = 0; I < 5; ++I) {
+ int Other;
+ // FIXME: We don't catch this because we use decl-context instead of scope.
+#pragma acc declare create(Local2, Arg2)
+ // OK, same scope.
+#pragma acc declare create(I, Other)
+// expected-error at +4 2{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -2 2{{previous reference is here}}
+// expected-error at +2 2{{variable referenced in 'copyin' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -4 2{{previous reference is here}}
+#pragma acc declare create(I, Other) copyin(I, Other)
+ }
+// expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(NonStaticMem)
+}
+
+template<typename T>
+struct DepStruct {
+ static const T DepStaticMem = 5;
+ static const int StaticMem = 5;
+ int NonStaticMem;
+// expected-error at +1{{no valid clauses specified in OpenACC 'declare' directive}}
+#pragma acc declare
+// expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(Global)
+ // OK, same scope.
+#pragma acc declare create(DepStaticMem)
+ // OK, same scope.
+#pragma acc declare create(StaticMem)
+// expected-error at +4{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -2{{previous reference is here}}
+// expected-error at +2{{variable referenced in 'copyin' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -4{{previous reference is here}}
+#pragma acc declare create(StaticMem) copyin(StaticMem)
+// expected-error at +4{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -9{{previous reference is here}}
+// expected-error at +2{{variable referenced in 'copyin' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -11{{previous reference is here}}
+#pragma acc declare create(DepStaticMem) copyin(DepStaticMem)
+
+ void Inline(int Arg) {
+// expected-error at +1{{no valid clauses specified in OpenACC 'declare' directive}}
+#pragma acc declare
+// expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(DepStaticMem)
+// expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(StaticMem)
+
+ T Local, Local2;
+ // OK, same scope.
+#pragma acc declare create(Local, Arg)
+// expected-error at +2 2{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -2 2{{previous reference is here}}
+#pragma acc declare create(Local, Local)
+
+ for (int I = 0; I < 5; ++I) {
+ int Other;
+ // FIXME: Since we approximate this as a decl-context, we can't check
+ // scope here.
+#pragma acc declare create(Local2)
+ // OK, same scope.
+#pragma acc declare create(I, Other)
+ // expected-error at +2 3{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
+ // expected-note at -2 3{{previous reference is here}}
+#pragma acc declare create(I, Other, I)
+ }
+ // expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(NonStaticMem)
+ }
+
+ void OutOfLine(int Arg);
+
+ template<typename U>
+ void TemplInline(U Arg, U Arg2) {
+// expected-error at +1{{no valid clauses specified in OpenACC 'declare' directive}}
+#pragma acc declare
+ // expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(DepStaticMem)
+ // expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(StaticMem)
+
+ T Local, Local2, Local3;
+ // OK, same scope.
+#pragma acc declare create(Local, Arg)
+// expected-error at +4{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -2{{previous reference is here}}
+// expected-error at +2{{variable referenced in 'present' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -4{{previous reference is here}}
+#pragma acc declare create(Local2, Arg) present(Local, Arg2)
+ {
+ // FIXME: We don't catch this, since we check decl-context not scopes.
+#pragma acc declare create(Local3)
+
+ // expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(NonStaticMem)
+ }
+ }
+ template<typename U>
+ void TemplOutline(U Arg);
+};
+
+template<typename T>
+void DepStruct<T>::OutOfLine(int Arg) {
+ // expected-error at +1{{no valid clauses specified in OpenACC 'declare' directive}}
+#pragma acc declare
+ // expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(StaticMem)
+
+ T Local, Local2;
+// OK, same scope.
+#pragma acc declare create(Local, Arg)
+
+ for (int I = 0; I < 5; ++I) {
+ int Other;
+ // FIXME: We don't catch this because we use decl-context instead of scope.
+#pragma acc declare create(Local2)
+ // OK, same scope.
+#pragma acc declare create(I, Other)
+ }
+ // expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(NonStaticMem)
+}
+template<typename T>
+template<typename U>
+void DepStruct<T>::TemplOutline(U Arg) {
+// expected-error at +1{{no valid clauses specified in OpenACC 'declare' directive}}
+#pragma acc declare
+ // expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(DepStaticMem)
+ // expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(StaticMem)
+
+ T Local, Local2;
+// OK, same scope.
+#pragma acc declare create(Local, Arg)
+
+ {
+ // FIXME: We could potentially fix this someday, but as we don't have
+ // 'scope' information like this during template instantiation, we have to
+ // permit this.
+#pragma acc declare create(Local2)
+ }
+ // expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(NonStaticMem)
+}
+
+void use() {
+ DepStruct<int> DS;
+ DS.Inline(1);
+ DS.OutOfLine(1);
+ DS.TemplInline(1, 2);
+ DS.TemplOutline(1);
+}
+
+// Only variable or array name.
+
+// expected-error at +1{{OpenACC variable on 'declare' construct is not a valid variable name or array name}}
+#pragma acc declare create(GlobalArray[0])
+// expected-error at +1{{OpenACC variable on 'declare' construct is not a valid variable name or array name}}
+#pragma acc declare create(GlobalArray[0: 1])
+
+struct S { int I; };
+// expected-error at +1{{OpenACC variable on 'declare' construct is not a valid variable name or array name}}
+#pragma acc declare create(S{}.I)
+
+int GS1, GS2, GS3, GS4, GS5, *GS6, GS7, GS8;
+
+// Only create, copyin, deviceptr, device-resident, link at NS scope.
+// expected-error at +3{{OpenACC 'copy' clause on a 'declare' directive is not allowed at global or namespace scope}}
+// expected-error at +2{{OpenACC 'copyout' clause on a 'declare' directive is not allowed at global or namespace scope}}
+// expected-error at +1{{OpenACC 'present' clause on a 'declare' directive is not allowed at global or namespace scope}}
+#pragma acc declare copy(GS1) copyin(GS2), copyout(GS3), create(GS4), present(GS5), deviceptr(GS6), device_resident(GS7), link(GS8)
+
+void ExternVar() {
+ extern int I, I2, I3, I4, I5, *I6, I7, I8;
+// expected-error at +3{{'extern' variable may not be referenced by 'copy' clause on an OpenACC 'declare' directive}}
+// expected-error at +2{{'extern' variable may not be referenced by 'copyout' clause on an OpenACC 'declare' directive}}
+// expected-error at +1{{'extern' variable may not be referenced by 'present' clause on an OpenACC 'declare' directive}}
+#pragma acc declare copy(I) copyin(I2), copyout(I3), create(I4), present(I5), deviceptr(I6), device_resident(I7), link(I8)
+}
+
+// Link can only have global, namespace, or extern vars.
+#pragma acc declare link(Global, GlobalArray)
+
+struct Struct2 {
+ static const int StaticMem = 5;
+ // expected-error at +1{{variable referenced by 'link' clause not in global or namespace scope must be marked 'extern'}}
+#pragma acc declare link(StaticMem)
+
+ void MemFunc(int I) {
+ int Local;
+ extern int ExternLocal;
+
+ // expected-error at +2{{variable referenced by 'link' clause not in global or namespace scope must be marked 'extern'}}
+ // expected-error at +1{{variable referenced by 'link' clause not in global or namespace scope must be marked 'extern'}}
+#pragma acc declare link(I, Local, ExternLocal)
+}
+};
+
diff --git a/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c b/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
index f56a1267fbad1..b4a705ba7d1c4 100644
--- a/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
+++ b/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
@@ -86,7 +86,7 @@ void uses() {
// expected-error at +1{{OpenACC 'deviceptr' clause is not valid on 'loop' directive}}
#pragma acc loop auto deviceptr(VarPtr)
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
+ // expected-error at +1{{OpenACC 'device_resident' clause is not valid on 'loop' directive}}
#pragma acc loop auto device_resident(VarPtr)
for(unsigned i = 0; i < 5; ++i);
// expected-error at +1{{OpenACC 'firstprivate' clause is not valid on 'loop' directive}}
@@ -95,7 +95,7 @@ void uses() {
// expected-error at +1{{OpenACC 'host' clause is not valid on 'loop' directive}}
#pragma acc loop auto host(Var)
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'link' not yet implemented}}
+ // expected-error at +1{{OpenACC 'link' clause is not valid on 'loop' directive}}
#pragma acc loop auto link(Var)
for(unsigned i = 0; i < 5; ++i);
// expected-error at +1{{OpenACC 'no_create' clause is not valid on 'loop' directive}}
@@ -220,7 +220,7 @@ void uses() {
// expected-error at +1{{OpenACC 'deviceptr' clause is not valid on 'loop' directive}}
#pragma acc loop deviceptr(VarPtr) auto
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
+ // expected-error at +1{{OpenACC 'device_resident' clause is not valid on 'loop' directive}}
#pragma acc loop device_resident(VarPtr) auto
for(unsigned i = 0; i < 5; ++i);
// expected-error at +1{{OpenACC 'firstprivate' clause is not valid on 'loop' directive}}
@@ -229,7 +229,7 @@ void uses() {
// expected-error at +1{{OpenACC 'host' clause is not valid on 'loop' directive}}
#pragma acc loop host(Var) auto
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'link' not yet implemented}}
+ // expected-error at +1{{OpenACC 'link' clause is not valid on 'loop' directive}}
#pragma acc loop link(Var) auto
for(unsigned i = 0; i < 5; ++i);
// expected-error at +1{{OpenACC 'no_create' clause is not valid on 'loop' directive}}
@@ -355,7 +355,7 @@ void uses() {
// expected-error at +1{{OpenACC 'deviceptr' clause is not valid on 'loop' directive}}
#pragma acc loop independent deviceptr(VarPtr)
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
+ // expected-error at +1{{OpenACC 'device_resident' clause is not valid on 'loop' directive}}
#pragma acc loop independent device_resident(VarPtr)
for(unsigned i = 0; i < 5; ++i);
// expected-error at +1{{OpenACC 'firstprivate' clause is not valid on 'loop' directive}}
@@ -364,7 +364,7 @@ void uses() {
// expected-error at +1{{OpenACC 'host' clause is not valid on 'loop' directive}}
#pragma acc loop independent host(Var)
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'link' not yet implemented}}
+ // expected-error at +1{{OpenACC 'link' clause is not valid on 'loop' directive}}
#pragma acc loop independent link(Var)
for(unsigned i = 0; i < 5; ++i);
// expected-error at +1{{OpenACC 'no_create' clause is not valid on 'loop' directive}}
@@ -489,7 +489,7 @@ void uses() {
// expected-error at +1{{OpenACC 'deviceptr' clause is not valid on 'loop' directive}}
#pragma acc loop deviceptr(VarPtr) independent
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
+ // expected-error at +1{{OpenACC 'device_resident' clause is not valid on 'loop' directive}}
#pragma acc loop device_resident(VarPtr) independent
for(unsigned i = 0; i < 5; ++i);
// expected-error at +1{{OpenACC 'firstprivate' clause is not valid on 'loop' directive}}
@@ -498,7 +498,7 @@ void uses() {
// expected-error at +1{{OpenACC 'host' clause is not valid on 'loop' directive}}
#pragma acc loop host(Var) independent
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'link' not yet implemented}}
+ // expected-error at +1{{OpenACC 'link' clause is not valid on 'loop' directive}}
#pragma acc loop link(Var) independent
for(unsigned i = 0; i < 5; ++i);
// expected-error at +1{{OpenACC 'no_create' clause is not valid on 'loop' directive}}
@@ -632,7 +632,7 @@ void uses() {
// expected-error at +1{{OpenACC 'deviceptr' clause is not valid on 'loop' directive}}
#pragma acc loop seq deviceptr(VarPtr)
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
+ // expected-error at +1{{OpenACC 'device_resident' clause is not valid on 'loop' directive}}
#pragma acc loop seq device_resident(VarPtr)
for(unsigned i = 0; i < 5; ++i);
// expected-error at +1{{OpenACC 'firstprivate' clause is not valid on 'loop' directive}}
@@ -641,7 +641,7 @@ void uses() {
// expected-error at +1{{OpenACC 'host' clause is not valid on 'loop' directive}}
#pragma acc loop seq host(Var)
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'link' not yet implemented}}
+ // expected-error at +1{{OpenACC 'link' clause is not valid on 'loop' directive}}
#pragma acc loop seq link(Var)
for(unsigned i = 0; i < 5; ++i);
// expected-error at +1{{OpenACC 'no_create' clause is not valid on 'loop' directive}}
@@ -772,7 +772,7 @@ void uses() {
// expected-error at +1{{OpenACC 'deviceptr' clause is not valid on 'loop' directive}}
#pragma acc loop deviceptr(VarPtr) seq
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
+ // expected-error at +1{{OpenACC 'device_resident' clause is not valid on 'loop' directive}}
#pragma acc loop device_resident(VarPtr) seq
for(unsigned i = 0; i < 5; ++i);
// expected-error at +1{{OpenACC 'firstprivate' clause is not valid on 'loop' directive}}
@@ -781,7 +781,7 @@ void uses() {
// expected-error at +1{{OpenACC 'host' clause is not valid on 'loop' directive}}
#pragma acc loop host(Var) seq
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'link' not yet implemented}}
+ // expected-error at +1{{OpenACC 'link' clause is not valid on 'loop' directive}}
#pragma acc loop link(Var) seq
for(unsigned i = 0; i < 5; ++i);
// expected-error at +1{{OpenACC 'no_create' clause is not valid on 'loop' directive}}
diff --git a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
index 2c1189bc647d0..f05ca11b143c0 100644
--- a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
@@ -98,8 +98,7 @@ void uses() {
// expected-error at +1{{OpenACC 'deviceptr' clause is not valid on 'loop' directive}}
#pragma acc loop device_type(*) deviceptr(VarPtr)
for(int i = 0; i < 5; ++i);
- // expected-error at +2{{OpenACC clause 'device_resident' may not follow a 'device_type' clause in a 'loop' construct}}
- // expected-note at +1{{previous clause is here}}
+ // expected-error at +1{{OpenACC 'device_resident' clause is not valid on 'loop' directive}}
#pragma acc loop device_type(*) device_resident(VarPtr)
for(int i = 0; i < 5; ++i);
// expected-error at +1{{OpenACC 'firstprivate' clause is not valid on 'loop' directive}}
@@ -108,8 +107,7 @@ void uses() {
// expected-error at +1{{OpenACC 'host' clause is not valid on 'loop' directive}}
#pragma acc loop device_type(*) host(Var)
for(int i = 0; i < 5; ++i);
- // expected-error at +2{{OpenACC clause 'link' may not follow a 'device_type' clause in a 'loop' construct}}
- // expected-note at +1{{previous clause is here}}
+ // expected-error at +1{{OpenACC 'link' clause is not valid on 'loop' directive}}
#pragma acc loop device_type(*) link(Var)
for(int i = 0; i < 5; ++i);
// expected-error at +1{{OpenACC 'no_create' clause is not valid on 'loop' directive}}
diff --git a/clang/test/SemaOpenACC/unimplemented-construct.c b/clang/test/SemaOpenACC/unimplemented-construct.c
index b22bfa5a69bb0..261c855cdd010 100644
--- a/clang/test/SemaOpenACC/unimplemented-construct.c
+++ b/clang/test/SemaOpenACC/unimplemented-construct.c
@@ -4,36 +4,36 @@
#pragma acc routine
struct S {
-// expected-warning at +1{{OpenACC construct 'declare' not yet implemented, pragma ignored}}
-#pragma acc declare
+// expected-warning at +1{{OpenACC construct 'routine' not yet implemented, pragma ignored}}
+#pragma acc routine
int foo;
};
void func() {
-// expected-warning at +1{{OpenACC construct 'declare' not yet implemented, pragma ignored}}
-#pragma acc declare
+// expected-warning at +1{{OpenACC construct 'routine' not yet implemented, pragma ignored}}
+#pragma acc routine
int foo;
-// expected-warning at +1{{OpenACC construct 'declare' not yet implemented, pragma ignored}}
-#pragma acc declare
+// expected-warning at +1{{OpenACC construct 'routine' not yet implemented, pragma ignored}}
+#pragma acc routine
{
-// expected-warning at +1{{OpenACC construct 'declare' not yet implemented, pragma ignored}}
-#pragma acc declare
+// expected-warning at +1{{OpenACC construct 'routine' not yet implemented, pragma ignored}}
+#pragma acc routine
{
-// expected-warning at +1{{OpenACC construct 'declare' not yet implemented, pragma ignored}}
-#pragma acc declare
+// expected-warning at +1{{OpenACC construct 'routine' not yet implemented, pragma ignored}}
+#pragma acc routine
}
}
-// expected-warning at +1{{OpenACC construct 'declare' not yet implemented, pragma ignored}}
-#pragma acc declare
+// expected-warning at +1{{OpenACC construct 'routine' not yet implemented, pragma ignored}}
+#pragma acc routine
while(0){}
-// expected-warning at +1{{OpenACC construct 'declare' not yet implemented, pragma ignored}}
-#pragma acc declare
+// expected-warning at +1{{OpenACC construct 'routine' not yet implemented, pragma ignored}}
+#pragma acc routine
for(;;){}
-// expected-warning at +1{{OpenACC construct 'declare' not yet implemented, pragma ignored}}
-#pragma acc declare
+// expected-warning at +1{{OpenACC construct 'routine' not yet implemented, pragma ignored}}
+#pragma acc routine
};
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index c710816bd24d0..479490442f3c8 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2905,6 +2905,13 @@ void OpenACCClauseEnqueue::VisitNoCreateClause(const OpenACCNoCreateClause &C) {
void OpenACCClauseEnqueue::VisitCopyClause(const OpenACCCopyClause &C) {
VisitVarList(C);
}
+void OpenACCClauseEnqueue::VisitLinkClause(const OpenACCLinkClause &C) {
+ VisitVarList(C);
+}
+void OpenACCClauseEnqueue::VisitDeviceResidentClause(
+ const OpenACCDeviceResidentClause &C) {
+ VisitVarList(C);
+}
void OpenACCClauseEnqueue::VisitCopyInClause(const OpenACCCopyInClause &C) {
VisitVarList(C);
}
@@ -7249,6 +7256,7 @@ CXCursor clang_getCursorDefinition(CXCursor C) {
case Decl::LifetimeExtendedTemporary:
case Decl::RequiresExprBody:
case Decl::UnresolvedUsingIfExists:
+ case Decl::OpenACCDeclare:
return C;
// Declaration kinds that don't make any sense here, but are
More information about the cfe-commits
mailing list