mirrors/llvm-project
0
0
Fork 0
mirror of https://github.com/llvm/llvm-project.git synced 2025-04-18 19:26:45 +00:00
Code Issues Projects Releases Wiki Activity

[OpenACC] Implement 'declare' construct AST/Sema

Browse Source 
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.
This commit is contained in:
erichkeane 2025-01-21 11:27:23 -08:00
parent ccf1bfc1d5
commit 5d7d66ba0d
50 changed files with 1788 additions and 169 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

106
clang/include/clang/AST/DeclOpenACC.h Normal file
View File

@ -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

1
clang/include/clang/AST/DeclVisitor.h
View File

@ -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"

2
clang/include/clang/AST/JSONNodeDumper.h
View File

@ -281,6 +281,8 @@ public:
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);

49
clang/include/clang/AST/OpenACCClause.h
View File

@ -1171,6 +1171,55 @@ public:
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); }

4
clang/include/clang/AST/RecursiveASTVisitor.h
View File

@ -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(

1
clang/include/clang/AST/TextNodeDumper.h
View File

@ -423,6 +423,7 @@ public:
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);

1
clang/include/clang/Basic/DeclNodes.td
View File

@ -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">;

24
clang/include/clang/Basic/DiagnosticSemaKinds.td
View File

@ -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">;

2
clang/include/clang/Basic/OpenACCClauses.def
View File

@ -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)

6
clang/include/clang/Parse/Parser.h
View File

@ -3756,9 +3756,11 @@ private:
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

1
clang/include/clang/Sema/SemaBase.h
View File

@ -42,6 +42,7 @@ public:
ASTContext &getASTContext() const;
DiagnosticsEngine &getDiagnostics() const;
const LangOptions &getLangOpts() const;
DeclContext *getCurContext() const;
/// Helper class that creates diagnostics with optional
/// template instantiation stacks.

48
clang/include/clang/Sema/SemaOpenACC.h
View File

@ -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 @@ private:
/// 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 @@ public:
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 @@ public:
}
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 @@ public:
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 @@ public:
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 @@ public:
/// 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 @@ public:
/// 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.

5
clang/include/clang/Serialization/ASTBitCodes.h
View File

@ -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.

3
clang/include/clang/Serialization/ASTRecordReader.h
View File

@ -278,7 +278,8 @@ public:
/// 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.

1
clang/lib/AST/ASTStructuralEquivalence.cpp
View File

@ -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"

1
clang/lib/AST/CMakeLists.txt
View File

@ -50,6 +50,7 @@ add_clang_library(clangAST
DeclFriend.cpp
DeclGroup.cpp
DeclObjC.cpp
DeclOpenACC.cpp
DeclOpenMP.cpp
DeclPrinter.cpp
DeclTemplate.cpp

2
clang/lib/AST/DeclBase.cpp
View File

@ -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;
}

33
clang/lib/AST/DeclOpenACC.cpp Normal file
View File

@ -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);
}

12
clang/lib/AST/DeclPrinter.cpp
View File

@ -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());
}
}

2
clang/lib/AST/JSONNodeDumper.cpp
View File

@ -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()));
}

36
clang/lib/AST/OpenACCClause.cpp
View File

@ -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())

6
clang/lib/AST/StmtPrinter.cpp
View File

@ -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) {

10
clang/lib/AST/StmtProfile.cpp
View File

@ -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);

13
clang/lib/AST/TextNodeDumper.cpp
View File

@ -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(); });

9
clang/lib/CodeGen/CGDecl.cpp
View File

@ -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);
}

3
clang/lib/CodeGen/CodeGenModule.h
View File

@ -1569,6 +1569,9 @@ public:
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);

36
clang/lib/Parse/ParseOpenACC.cpp
View File

@ -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.

1
clang/lib/Sema/SemaBase.cpp
View File

@ -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.

93
clang/lib/Sema/SemaOpenACC.cpp
View File

@ -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) {

254
clang/lib/Sema/SemaOpenACCClause.cpp
View File

@ -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;
}

228
clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
View File

@ -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

18
clang/lib/Sema/TreeTransform.h
View File

@ -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) {

1
clang/lib/Serialization/ASTCommon.cpp
View File

@ -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

16
clang/lib/Serialization/ASTReader.cpp
View File

@ -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");

13
clang/lib/Serialization/ASTReaderDecl.cpp
View File

@ -414,6 +414,8 @@ public:
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");

15
clang/lib/Serialization/ASTWriter.cpp
View File

@ -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");

12
clang/lib/Serialization/ASTWriterDecl.cpp
View File

@ -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
//===----------------------------------------------------------------------===//

52
clang/test/AST/ast-print-openacc-declare-construct.cpp Normal file
View File

@ -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)
}

30
clang/test/ParserOpenACC/parse-clauses.c
View File

@ -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@+2{{expected ','}}
// expected-warning@+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@+3{{expected ','}}
// expected-error@+2{{OpenACC variable on 'declare' construct is not a valid variable name or array name}}
// expected-error@+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@+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@+2{{expected ','}}
// expected-warning@+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@+2{{OpenACC variable on 'declare' construct is not a valid variable name or array name}}
// expected-error@+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@+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@+3{{expected ','}}
// expected-error@+2{{OpenACC variable on 'declare' construct is not a valid variable name or array name}}
// expected-error@+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@+2{{OpenACC variable on 'declare' construct is not a valid variable name or array name}}
// expected-error@+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@+1{{expected ','}}
#pragma acc update host(s.array[s.value] s.array[s.value :5] )

2
clang/test/ParserOpenACC/parse-constructs.c
View File

@ -131,7 +131,7 @@ void func() {
// expected-error@+2{{invalid OpenACC clause 'clause'}}
// expected-warning@+1{{OpenACC construct 'declare' not yet implemented, pragma ignored}}
// expected-error@+1{{no valid clauses specified in OpenACC 'declare' directive}}
#pragma acc declare clause list
for(;;){}
// expected-error@+1{{invalid OpenACC clause 'clause'}}

24
clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
View File

@ -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@+1{{OpenACC clause 'device_resident' not yet implemented}}
// expected-error@+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@+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@+1{{OpenACC clause 'link' not yet implemented}}
// expected-error@+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@+1{{OpenACC clause 'device_resident' not yet implemented}}
// expected-error@+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@+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@+1{{OpenACC clause 'link' not yet implemented}}
// expected-error@+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@+1{{OpenACC clause 'device_resident' not yet implemented}}
// expected-error@+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@+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@+1{{OpenACC clause 'link' not yet implemented}}
// expected-error@+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@+1{{OpenACC clause 'device_resident' not yet implemented}}
// expected-error@+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@+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@+1{{OpenACC clause 'link' not yet implemented}}
// expected-error@+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@+1{{OpenACC clause 'device_resident' not yet implemented}}
// expected-error@+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@+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@+1{{OpenACC clause 'link' not yet implemented}}
// expected-error@+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@+1{{OpenACC clause 'device_resident' not yet implemented}}
// expected-error@+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@+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@+1{{OpenACC clause 'link' not yet implemented}}
// expected-error@+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

6
clang/test/SemaOpenACC/combined-construct-device_type-clause.c
View File

@ -107,8 +107,7 @@ void uses() {
// expected-note@+1{{previous clause is here}}
#pragma acc serial loop device_type(*) deviceptr(VarPtr)
for(int i = 0; i < 5; ++i);
// expected-error@+2{{OpenACC clause 'device_resident' may not follow a 'device_type' clause in a 'kernels loop' construct}}
// expected-note@+1{{previous clause is here}}
// expected-error@+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@+2{{OpenACC clause 'firstprivate' may not follow a 'device_type' clause in a 'parallel loop' construct}}
@ -118,8 +117,7 @@ void uses() {
// expected-error@+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@+2{{OpenACC clause 'link' may not follow a 'device_type' clause in a 'parallel loop' construct}}
// expected-note@+1{{previous clause is here}}
// expected-error@+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@+2{{OpenACC clause 'no_create' may not follow a 'device_type' clause in a 'serial loop' construct}}

6
clang/test/SemaOpenACC/compute-construct-device_type-clause.c
View File

@ -113,8 +113,7 @@ void uses() {
// expected-note@+1{{previous clause is here}}
#pragma acc kernels device_type(*) deviceptr(VarPtr)
while(1);
// expected-error@+2{{OpenACC clause 'device_resident' may not follow a 'device_type' clause in a 'kernels' construct}}
// expected-note@+1{{previous clause is here}}
// expected-error@+1{{OpenACC 'device_resident' clause is not valid on 'kernels' directive}}
#pragma acc kernels device_type(*) device_resident(VarPtr)
while(1);
// expected-error@+2{{OpenACC clause 'firstprivate' may not follow a 'device_type' clause in a 'parallel' construct}}
@ -124,8 +123,7 @@ void uses() {
// expected-error@+1{{OpenACC 'host' clause is not valid on 'kernels' directive}}
#pragma acc kernels device_type(*) host(Var)
while(1);
// expected-error@+2{{OpenACC clause 'link' may not follow a 'device_type' clause in a 'kernels' construct}}
// expected-note@+1{{previous clause is here}}
// expected-error@+1{{OpenACC 'link' clause is not valid on 'kernels' directive}}
#pragma acc kernels device_type(*) link(Var)
while(1);
// expected-error@+2{{OpenACC clause 'no_create' may not follow a 'device_type' clause in a 'kernels' construct}}

1
clang/test/SemaOpenACC/data-construct.cpp
View File

@ -165,7 +165,6 @@ struct HasMembers {
// expected-error@+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@+1{{OpenACC variable in 'use_device' clause is not a valid variable name or array name}}
#pragma acc host_data use_device(Member)
;
}

390
clang/test/SemaOpenACC/declare-construct-ast.cpp Normal file
View File

@ -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

308
clang/test/SemaOpenACC/declare-construct.cpp Normal file
View File

@ -0,0 +1,308 @@
// RUN: %clang_cc1 %s -fopenacc -verify
int *Global;
int GlobalArray[5];
// expected-error@+1{{no valid clauses specified in OpenACC 'declare' directive}}
#pragma acc declare
namespace NS {
int *NSVar;
int NSArray[5];
// expected-error@+2{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
// expected-error@+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@+4{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
// expected-note@-3{{previous reference is here}}
// expected-error@+2{{variable referenced in 'copyin' clause of OpenACC 'declare' directive was already referenced}}
// expected-note@-5{{previous reference is here}}
#pragma acc declare create(NSVar) copyin(NSVar)
// expected-error@+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@+3{{OpenACC 'copy' clause on a 'declare' directive is not allowed at global or namespace scope}}
// expected-error@+2{{OpenACC 'copyout' clause on a 'declare' directive is not allowed at global or namespace scope}}
// expected-error@+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@+3{{OpenACC 'copy' clause on a 'declare' directive is not allowed at global or namespace scope}}
// expected-error@+2{{OpenACC 'copyout' clause on a 'declare' directive is not allowed at global or namespace scope}}
// expected-error@+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@+2{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
// expected-error@+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@+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@+4{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
// expected-note@-2{{previous reference is here}}
// expected-error@+2{{variable referenced in 'copyin' clause of OpenACC 'declare' directive was already referenced}}
// expected-note@-4{{previous reference is here}}
#pragma acc declare create(StaticMem) copyin(StaticMem)
// expected-error@+1{{no valid clauses specified in OpenACC 'declare' directive}}
#pragma acc declare
void Inline(int Arg) {
// expected-error@+1{{no valid clauses specified in OpenACC 'declare' directive}}
#pragma acc declare
// expected-error@+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@+2{{variable referenced in 'copyin' clause of OpenACC 'declare' directive was already referenced}}
// expected-note@+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@+4 2{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
// expected-note@-2 2{{previous reference is here}}
// expected-error@+2 2{{variable referenced in 'copyin' clause of OpenACC 'declare' directive was already referenced}}
// expected-note@-4 2{{previous reference is here}}
#pragma acc declare create(I, Other) copyin(I, Other)
// expected-error@+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@+1{{no valid clauses specified in OpenACC 'declare' directive}}
#pragma acc declare
// expected-error@+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@+4{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
// expected-note@-2{{previous reference is here}}
// expected-error@+2{{variable referenced in 'copyin' clause of OpenACC 'declare' directive was already referenced}}
// expected-note@-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@+4 2{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
// expected-note@-2 2{{previous reference is here}}
// expected-error@+2 2{{variable referenced in 'copyin' clause of OpenACC 'declare' directive was already referenced}}
// expected-note@-4 2{{previous reference is here}}
#pragma acc declare create(I, Other) copyin(I, Other)
}
// expected-error@+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@+1{{no valid clauses specified in OpenACC 'declare' directive}}
#pragma acc declare
// expected-error@+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@+4{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
// expected-note@-2{{previous reference is here}}
// expected-error@+2{{variable referenced in 'copyin' clause of OpenACC 'declare' directive was already referenced}}
// expected-note@-4{{previous reference is here}}
#pragma acc declare create(StaticMem) copyin(StaticMem)
// expected-error@+4{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
// expected-note@-9{{previous reference is here}}
// expected-error@+2{{variable referenced in 'copyin' clause of OpenACC 'declare' directive was already referenced}}
// expected-note@-11{{previous reference is here}}
#pragma acc declare create(DepStaticMem) copyin(DepStaticMem)
void Inline(int Arg) {
// expected-error@+1{{no valid clauses specified in OpenACC 'declare' directive}}
#pragma acc declare
// expected-error@+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@+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@+2 2{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
// expected-note@-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@+2 3{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
// expected-note@-2 3{{previous reference is here}}
#pragma acc declare create(I, Other, I)
}
// expected-error@+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@+1{{no valid clauses specified in OpenACC 'declare' directive}}
#pragma acc declare
// expected-error@+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@+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@+4{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
// expected-note@-2{{previous reference is here}}
// expected-error@+2{{variable referenced in 'present' clause of OpenACC 'declare' directive was already referenced}}
// expected-note@-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@+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@+1{{no valid clauses specified in OpenACC 'declare' directive}}
#pragma acc declare
// expected-error@+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@+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@+1{{no valid clauses specified in OpenACC 'declare' directive}}
#pragma acc declare
// expected-error@+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@+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@+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@+1{{OpenACC variable on 'declare' construct is not a valid variable name or array name}}
#pragma acc declare create(GlobalArray[0])
// expected-error@+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@+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@+3{{OpenACC 'copy' clause on a 'declare' directive is not allowed at global or namespace scope}}
// expected-error@+2{{OpenACC 'copyout' clause on a 'declare' directive is not allowed at global or namespace scope}}
// expected-error@+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@+3{{'extern' variable may not be referenced by 'copy' clause on an OpenACC 'declare' directive}}
// expected-error@+2{{'extern' variable may not be referenced by 'copyout' clause on an OpenACC 'declare' directive}}
// expected-error@+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@+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@+2{{variable referenced by 'link' clause not in global or namespace scope must be marked 'extern'}}
// expected-error@+1{{variable referenced by 'link' clause not in global or namespace scope must be marked 'extern'}}
#pragma acc declare link(I, Local, ExternLocal)
}
};

24
clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
View File

@ -86,7 +86,7 @@ void uses() {
// expected-error@+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@+1{{OpenACC clause 'device_resident' not yet implemented}}
// expected-error@+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@+1{{OpenACC 'firstprivate' clause is not valid on 'loop' directive}}
@ -95,7 +95,7 @@ void uses() {
// expected-error@+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@+1{{OpenACC clause 'link' not yet implemented}}
// expected-error@+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@+1{{OpenACC 'no_create' clause is not valid on 'loop' directive}}
@ -220,7 +220,7 @@ void uses() {
// expected-error@+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@+1{{OpenACC clause 'device_resident' not yet implemented}}
// expected-error@+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@+1{{OpenACC 'firstprivate' clause is not valid on 'loop' directive}}
@ -229,7 +229,7 @@ void uses() {
// expected-error@+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@+1{{OpenACC clause 'link' not yet implemented}}
// expected-error@+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@+1{{OpenACC 'no_create' clause is not valid on 'loop' directive}}
@ -355,7 +355,7 @@ void uses() {
// expected-error@+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@+1{{OpenACC clause 'device_resident' not yet implemented}}
// expected-error@+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@+1{{OpenACC 'firstprivate' clause is not valid on 'loop' directive}}
@ -364,7 +364,7 @@ void uses() {
// expected-error@+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@+1{{OpenACC clause 'link' not yet implemented}}
// expected-error@+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@+1{{OpenACC 'no_create' clause is not valid on 'loop' directive}}
@ -489,7 +489,7 @@ void uses() {
// expected-error@+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@+1{{OpenACC clause 'device_resident' not yet implemented}}
// expected-error@+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@+1{{OpenACC 'firstprivate' clause is not valid on 'loop' directive}}
@ -498,7 +498,7 @@ void uses() {
// expected-error@+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@+1{{OpenACC clause 'link' not yet implemented}}
// expected-error@+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@+1{{OpenACC 'no_create' clause is not valid on 'loop' directive}}
@ -632,7 +632,7 @@ void uses() {
// expected-error@+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@+1{{OpenACC clause 'device_resident' not yet implemented}}
// expected-error@+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@+1{{OpenACC 'firstprivate' clause is not valid on 'loop' directive}}
@ -641,7 +641,7 @@ void uses() {
// expected-error@+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@+1{{OpenACC clause 'link' not yet implemented}}
// expected-error@+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@+1{{OpenACC 'no_create' clause is not valid on 'loop' directive}}
@ -772,7 +772,7 @@ void uses() {
// expected-error@+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@+1{{OpenACC clause 'device_resident' not yet implemented}}
// expected-error@+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@+1{{OpenACC 'firstprivate' clause is not valid on 'loop' directive}}
@ -781,7 +781,7 @@ void uses() {
// expected-error@+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@+1{{OpenACC clause 'link' not yet implemented}}
// expected-error@+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@+1{{OpenACC 'no_create' clause is not valid on 'loop' directive}}

6
clang/test/SemaOpenACC/loop-construct-device_type-clause.c
View File

@ -98,8 +98,7 @@ void uses() {
// expected-error@+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@+2{{OpenACC clause 'device_resident' may not follow a 'device_type' clause in a 'loop' construct}}
// expected-note@+1{{previous clause is here}}
// expected-error@+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@+1{{OpenACC 'firstprivate' clause is not valid on 'loop' directive}}
@ -108,8 +107,7 @@ void uses() {
// expected-error@+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@+2{{OpenACC clause 'link' may not follow a 'device_type' clause in a 'loop' construct}}
// expected-note@+1{{previous clause is here}}
// expected-error@+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@+1{{OpenACC 'no_create' clause is not valid on 'loop' directive}}

32
clang/test/SemaOpenACC/unimplemented-construct.c
View File

@ -4,36 +4,36 @@
#pragma acc routine
struct S {
// expected-warning@+1{{OpenACC construct 'declare' not yet implemented, pragma ignored}}
#pragma acc declare
// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}}
#pragma acc routine
int foo;
};
void func() {
// expected-warning@+1{{OpenACC construct 'declare' not yet implemented, pragma ignored}}
#pragma acc declare
// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}}
#pragma acc routine
int foo;
// expected-warning@+1{{OpenACC construct 'declare' not yet implemented, pragma ignored}}
#pragma acc declare
// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}}
#pragma acc routine
{
// expected-warning@+1{{OpenACC construct 'declare' not yet implemented, pragma ignored}}
#pragma acc declare
// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}}
#pragma acc routine
{
// expected-warning@+1{{OpenACC construct 'declare' not yet implemented, pragma ignored}}
#pragma acc declare
// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}}
#pragma acc routine
}
}
// expected-warning@+1{{OpenACC construct 'declare' not yet implemented, pragma ignored}}
#pragma acc declare
// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}}
#pragma acc routine
while(0){}
// expected-warning@+1{{OpenACC construct 'declare' not yet implemented, pragma ignored}}
#pragma acc declare
// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}}
#pragma acc routine
for(;;){}
// expected-warning@+1{{OpenACC construct 'declare' not yet implemented, pragma ignored}}
#pragma acc declare
// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}}
#pragma acc routine
};

8
clang/tools/libclang/CIndex.cpp
View File

@ -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