[clang] [OpenACC] Loop construct basic Sema and AST work (PR #93742)
Erich Keane via cfe-commits
cfe-commits at lists.llvm.org
Thu May 30 08:02:16 PDT 2024
https://github.com/erichkeane updated https://github.com/llvm/llvm-project/pull/93742
>From 51792c63e8fee3ad662174d6a53bbca195cbc1b4 Mon Sep 17 00:00:00 2001
From: erichkeane <ekeane at nvidia.com>
Date: Fri, 24 May 2024 12:12:29 -0700
Subject: [PATCH 1/3] [OpenACC] Loop construct basic Sema and AST work
This patch implements the 'loop' construct AST, as well as the basic
appertainment rule. Additionally, it sets up the 'parent' compute
construct, which is necessary for codegen/other diagnostics.
A 'loop' can apply to a for or range-for loop, otherwise it has no other
restrictions (though some of its clauses do).
---
clang/include/clang-c/Index.h | 6 +-
clang/include/clang/AST/RecursiveASTVisitor.h | 2 +
clang/include/clang/AST/StmtOpenACC.h | 71 ++-
clang/include/clang/AST/TextNodeDumper.h | 1 +
.../clang/Basic/DiagnosticSemaKinds.td | 3 +
clang/include/clang/Basic/StmtNodes.td | 1 +
clang/include/clang/Sema/SemaOpenACC.h | 27 +-
.../include/clang/Serialization/ASTBitCodes.h | 1 +
clang/lib/AST/StmtOpenACC.cpp | 91 +++-
clang/lib/AST/StmtPrinter.cpp | 13 +
clang/lib/AST/StmtProfile.cpp | 8 +
clang/lib/AST/TextNodeDumper.cpp | 7 +
clang/lib/CodeGen/CGStmt.cpp | 3 +
clang/lib/CodeGen/CodeGenFunction.h | 7 +
clang/lib/Parse/ParseOpenACC.cpp | 8 +-
clang/lib/Sema/SemaExceptionSpec.cpp | 1 +
clang/lib/Sema/SemaOpenACC.cpp | 60 ++-
clang/lib/Sema/TreeTransform.h | 44 +-
clang/lib/Serialization/ASTReaderStmt.cpp | 11 +
clang/lib/Serialization/ASTWriterStmt.cpp | 6 +
clang/lib/StaticAnalyzer/Core/ExprEngine.cpp | 1 +
.../AST/ast-print-openacc-loop-construct.cpp | 9 +
clang/test/ParserOpenACC/parse-clauses.c | 410 ++++++++----------
clang/test/ParserOpenACC/parse-clauses.cpp | 6 +-
clang/test/ParserOpenACC/parse-constructs.c | 3 +-
.../compute-construct-default-clause.c | 3 +-
clang/test/SemaOpenACC/loop-ast.cpp | 164 +++++++
clang/test/SemaOpenACC/loop-loc-and-stmt.c | 38 ++
clang/test/SemaOpenACC/loop-loc-and-stmt.cpp | 80 ++++
clang/tools/libclang/CIndex.cpp | 9 +
clang/tools/libclang/CXCursor.cpp | 3 +
31 files changed, 843 insertions(+), 254 deletions(-)
create mode 100644 clang/test/AST/ast-print-openacc-loop-construct.cpp
create mode 100644 clang/test/SemaOpenACC/loop-ast.cpp
create mode 100644 clang/test/SemaOpenACC/loop-loc-and-stmt.c
create mode 100644 clang/test/SemaOpenACC/loop-loc-and-stmt.cpp
diff --git a/clang/include/clang-c/Index.h b/clang/include/clang-c/Index.h
index 365b607c74117..ce2282937f86c 100644
--- a/clang/include/clang-c/Index.h
+++ b/clang/include/clang-c/Index.h
@@ -2150,7 +2150,11 @@ enum CXCursorKind {
*/
CXCursor_OpenACCComputeConstruct = 320,
- CXCursor_LastStmt = CXCursor_OpenACCComputeConstruct,
+ /** OpenACC Loop Construct.
+ */
+ CXCursor_OpenACCLoopConstruct = 321,
+
+ CXCursor_LastStmt = CXCursor_OpenACCLoopConstruct,
/**
* Cursor that represents the translation unit itself.
diff --git a/clang/include/clang/AST/RecursiveASTVisitor.h b/clang/include/clang/AST/RecursiveASTVisitor.h
index 4bbb4380cdd7f..bb6ff334dd32d 100644
--- a/clang/include/clang/AST/RecursiveASTVisitor.h
+++ b/clang/include/clang/AST/RecursiveASTVisitor.h
@@ -3996,6 +3996,8 @@ bool RecursiveASTVisitor<Derived>::VisitOpenACCClauseList(
DEF_TRAVERSE_STMT(OpenACCComputeConstruct,
{ TRY_TO(TraverseOpenACCAssociatedStmtConstruct(S)); })
+DEF_TRAVERSE_STMT(OpenACCLoopConstruct,
+ { TRY_TO(TraverseOpenACCAssociatedStmtConstruct(S)); })
// FIXME: look at the following tricky-seeming exprs to see if we
// need to recurse on anything. These are ones that have methods
diff --git a/clang/include/clang/AST/StmtOpenACC.h b/clang/include/clang/AST/StmtOpenACC.h
index 04daf511f5871..b3aea09be03dd 100644
--- a/clang/include/clang/AST/StmtOpenACC.h
+++ b/clang/include/clang/AST/StmtOpenACC.h
@@ -113,6 +113,8 @@ class OpenACCAssociatedStmtConstruct : public OpenACCConstructStmt {
return const_cast<OpenACCAssociatedStmtConstruct *>(this)->children();
}
};
+
+class OpenACCLoopConstruct;
/// This class represents a compute construct, representing a 'Kind' of
/// `parallel', 'serial', or 'kernel'. These constructs are associated with a
/// 'structured block', defined as:
@@ -165,6 +167,11 @@ class OpenACCComputeConstruct final
}
void setStructuredBlock(Stmt *S) { setAssociatedStmt(S); }
+ // Serialization helper function that searches the structured block for 'loop'
+ // constructs that should be associated with this, and sets their parent
+ // compute construct to this one. This isn't necessary normally, since we have
+ // the ability to record the state during parsing.
+ void findAndSetChildLoops();
public:
static bool classof(const Stmt *T) {
@@ -176,12 +183,74 @@ class OpenACCComputeConstruct final
static OpenACCComputeConstruct *
Create(const ASTContext &C, OpenACCDirectiveKind K, SourceLocation BeginLoc,
SourceLocation DirectiveLoc, SourceLocation EndLoc,
- ArrayRef<const OpenACCClause *> Clauses, Stmt *StructuredBlock);
+ ArrayRef<const OpenACCClause *> Clauses, Stmt *StructuredBlock,
+ ArrayRef<OpenACCLoopConstruct *> AssociatedLoopConstructs);
Stmt *getStructuredBlock() { return getAssociatedStmt(); }
const Stmt *getStructuredBlock() const {
return const_cast<OpenACCComputeConstruct *>(this)->getStructuredBlock();
}
};
+/// This class represents a 'loop' construct. The 'loop' construct applies to a
+/// 'for' loop (or range-for loop), and is optionally associated with a Compute
+/// Construct.
+class OpenACCLoopConstruct final
+ : public OpenACCAssociatedStmtConstruct,
+ public llvm::TrailingObjects<OpenACCLoopConstruct,
+ const OpenACCClause *> {
+ // The compute construct this loop is associated with, or nullptr if this is
+ // an orphaned loop construct, or if it hasn't been set yet. Because we
+ // construct the directives at the end of their statement, the 'parent'
+ // construct is not yet available at the time of construction, so this needs
+ // to be set 'later'.
+ const OpenACCComputeConstruct *ParentComputeConstruct = nullptr;
+
+ friend class ASTStmtWriter;
+ friend class ASTStmtReader;
+ friend class ASTContext;
+ friend class OpenACCComputeConstruct;
+
+ OpenACCLoopConstruct(unsigned NumClauses);
+
+ OpenACCLoopConstruct(SourceLocation Start, SourceLocation DirLoc,
+ SourceLocation End,
+ ArrayRef<const OpenACCClause *> Clauses, Stmt *Loop);
+ void setLoop(Stmt *Loop);
+
+ void setParentComputeConstruct(OpenACCComputeConstruct *CC) {
+ assert(!ParentComputeConstruct && "Parent already set?");
+ ParentComputeConstruct = CC;
+ }
+
+public:
+ static bool classof(const Stmt *T) {
+ return T->getStmtClass() == OpenACCLoopConstructClass;
+ }
+
+ static OpenACCLoopConstruct *CreateEmpty(const ASTContext &C,
+ unsigned NumClauses);
+
+ static OpenACCLoopConstruct *
+ Create(const ASTContext &C, SourceLocation BeginLoc, SourceLocation DirLoc,
+ SourceLocation EndLoc, ArrayRef<const OpenACCClause *> Clauses,
+ Stmt *Loop);
+
+ Stmt *getLoop() { return getAssociatedStmt(); }
+ const Stmt *getLoop() const {
+ return const_cast<OpenACCLoopConstruct *>(this)->getLoop();
+ }
+
+ /// OpenACC 3.3 2.9:
+ /// An orphaned loop construct is a loop construct that is not lexically
+ /// enclosed within a compute construct. The parent compute construct of a
+ /// loop construct is the nearest compute construct that lexically contains
+ /// the loop construct.
+ bool isOrphanedLoopConstruct() const {
+ return ParentComputeConstruct == nullptr;
+ }
+ const OpenACCComputeConstruct *getParentComputeConstruct() const {
+ return ParentComputeConstruct;
+ }
+};
} // namespace clang
#endif // LLVM_CLANG_AST_STMTOPENACC_H
diff --git a/clang/include/clang/AST/TextNodeDumper.h b/clang/include/clang/AST/TextNodeDumper.h
index 63fa16c9ec47c..de82b438b0eb0 100644
--- a/clang/include/clang/AST/TextNodeDumper.h
+++ b/clang/include/clang/AST/TextNodeDumper.h
@@ -407,6 +407,7 @@ class TextNodeDumper
VisitLifetimeExtendedTemporaryDecl(const LifetimeExtendedTemporaryDecl *D);
void VisitHLSLBufferDecl(const HLSLBufferDecl *D);
void VisitOpenACCConstructStmt(const OpenACCConstructStmt *S);
+ void VisitOpenACCLoopConstruct(const OpenACCLoopConstruct *S);
};
} // namespace clang
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index f15cba63624ea..b3d985c5ff48c 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12406,6 +12406,9 @@ def err_acc_reduction_composite_type
def err_acc_reduction_composite_member_type :Error<
"OpenACC 'reduction' composite variable must not have non-scalar field">;
def note_acc_reduction_composite_member_loc : Note<"invalid field is here">;
+def err_acc_loop_not_for_loop
+ : Error<"OpenACC 'loop' construct can only be applied to a 'for' loop">;
+def note_acc_construct_here : Note<"'%0' construct is here">;
// AMDGCN builtins diagnostics
def err_amdgcn_global_load_lds_size_invalid_value : Error<"invalid size value">;
diff --git a/clang/include/clang/Basic/StmtNodes.td b/clang/include/clang/Basic/StmtNodes.td
index 305f19daa4a92..6ca08abdb14f0 100644
--- a/clang/include/clang/Basic/StmtNodes.td
+++ b/clang/include/clang/Basic/StmtNodes.td
@@ -302,3 +302,4 @@ def OpenACCConstructStmt : StmtNode<Stmt, /*abstract=*/1>;
def OpenACCAssociatedStmtConstruct
: StmtNode<OpenACCConstructStmt, /*abstract=*/1>;
def OpenACCComputeConstruct : StmtNode<OpenACCAssociatedStmtConstruct>;
+def OpenACCLoopConstruct : StmtNode<OpenACCAssociatedStmtConstruct>;
diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index 66144de4340a8..a5f2a8bf74657 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -15,6 +15,7 @@
#define LLVM_CLANG_SEMA_SEMAOPENACC_H
#include "clang/AST/DeclGroup.h"
+#include "clang/AST/StmtOpenACC.h"
#include "clang/Basic/OpenACCKinds.h"
#include "clang/Basic/SourceLocation.h"
#include "clang/Sema/Ownership.h"
@@ -25,6 +26,15 @@ namespace clang {
class OpenACCClause;
class SemaOpenACC : public SemaBase {
+private:
+ /// A collection of loop constructs in the compute construct scope that
+ /// haven't had their 'parent' compute construct set yet. Entires will only be
+ /// made to this list in the case where we know the loop isn't an orphan.
+ llvm::SmallVector<OpenACCLoopConstruct *> ParentlessLoopConstructs;
+ /// Whether we are inside of a compute construct, and should add loops to the
+ /// above collection.
+ bool InsideComputeConstruct = false;
+
public:
// Redeclaration of the version in OpenACCClause.h.
using DeviceTypeArgument = std::pair<IdentifierInfo *, SourceLocation>;
@@ -394,7 +404,8 @@ class SemaOpenACC : public SemaBase {
bool ActOnStartDeclDirective(OpenACCDirectiveKind K, SourceLocation StartLoc);
/// Called when we encounter an associated statement for our construct, this
/// should check legality of the statement as it appertains to this Construct.
- StmtResult ActOnAssociatedStmt(OpenACCDirectiveKind K, StmtResult AssocStmt);
+ StmtResult ActOnAssociatedStmt(SourceLocation DirectiveLoc,
+ OpenACCDirectiveKind K, StmtResult AssocStmt);
/// Called after the directive has been completely parsed, including the
/// declaration group or associated statement.
@@ -431,6 +442,20 @@ class SemaOpenACC : public SemaBase {
Expr *LowerBound,
SourceLocation ColonLocFirst, Expr *Length,
SourceLocation RBLoc);
+
+ /// Helper type for the registration/assignment of constructs that need to
+ /// 'know' about their parent constructs and hold a reference to them, such as
+ /// Loop needing its parent construct.
+ class AssociatedStmtRAII {
+ SemaOpenACC &SemaRef;
+ bool WasInsideComputeConstruct;
+ OpenACCDirectiveKind DirKind;
+ llvm::SmallVector<OpenACCLoopConstruct *> ParentlessLoopConstructs;
+
+ public:
+ AssociatedStmtRAII(SemaOpenACC &, OpenACCDirectiveKind);
+ ~AssociatedStmtRAII();
+ };
};
} // namespace clang
diff --git a/clang/include/clang/Serialization/ASTBitCodes.h b/clang/include/clang/Serialization/ASTBitCodes.h
index fe1bd47348be1..f59ff6af4c764 100644
--- a/clang/include/clang/Serialization/ASTBitCodes.h
+++ b/clang/include/clang/Serialization/ASTBitCodes.h
@@ -1946,6 +1946,7 @@ enum StmtCode {
// OpenACC Constructs
STMT_OPENACC_COMPUTE_CONSTRUCT,
+ STMT_OPENACC_LOOP_CONSTRUCT,
};
/// The kinds of designators that can occur in a
diff --git a/clang/lib/AST/StmtOpenACC.cpp b/clang/lib/AST/StmtOpenACC.cpp
index 47899b344c97a..e235449eb9959 100644
--- a/clang/lib/AST/StmtOpenACC.cpp
+++ b/clang/lib/AST/StmtOpenACC.cpp
@@ -12,6 +12,8 @@
#include "clang/AST/StmtOpenACC.h"
#include "clang/AST/ASTContext.h"
+#include "clang/AST/RecursiveASTVisitor.h"
+#include "clang/AST/StmtCXX.h"
using namespace clang;
OpenACCComputeConstruct *
@@ -26,11 +28,98 @@ OpenACCComputeConstruct::CreateEmpty(const ASTContext &C, unsigned NumClauses) {
OpenACCComputeConstruct *OpenACCComputeConstruct::Create(
const ASTContext &C, OpenACCDirectiveKind K, SourceLocation BeginLoc,
SourceLocation DirLoc, SourceLocation EndLoc,
- ArrayRef<const OpenACCClause *> Clauses, Stmt *StructuredBlock) {
+ ArrayRef<const OpenACCClause *> Clauses, Stmt *StructuredBlock,
+ ArrayRef<OpenACCLoopConstruct *> AssociatedLoopConstructs) {
void *Mem = C.Allocate(
OpenACCComputeConstruct::totalSizeToAlloc<const OpenACCClause *>(
Clauses.size()));
auto *Inst = new (Mem) OpenACCComputeConstruct(K, BeginLoc, DirLoc, EndLoc,
Clauses, StructuredBlock);
+
+ llvm::for_each(AssociatedLoopConstructs, [&](OpenACCLoopConstruct *C) {
+ C->setParentComputeConstruct(Inst);
+ });
+
+ return Inst;
+}
+
+void OpenACCComputeConstruct::findAndSetChildLoops() {
+ struct LoopConstructFinder : RecursiveASTVisitor<LoopConstructFinder> {
+ OpenACCComputeConstruct *Construct = nullptr;
+
+ LoopConstructFinder(OpenACCComputeConstruct *Construct)
+ : Construct(Construct) {}
+
+ bool VisitOpenACCComputeConstruct(OpenACCComputeConstruct *C) {
+ // Stop searching if we find a compute construct.
+ return false;
+ }
+ bool VisitOpenACCLoopConstruct(OpenACCLoopConstruct *C) {
+ // Stop searching if we find a loop construct, after taking ownership of
+ // it.
+ C->setParentComputeConstruct(Construct);
+ return false;
+ }
+ };
+
+ LoopConstructFinder f(this);
+ f.TraverseStmt(getAssociatedStmt());
+}
+
+OpenACCLoopConstruct::OpenACCLoopConstruct(unsigned NumClauses)
+ : OpenACCAssociatedStmtConstruct(
+ OpenACCLoopConstructClass, OpenACCDirectiveKind::Loop,
+ SourceLocation{}, SourceLocation{}, SourceLocation{},
+ /*AssociatedStmt=*/nullptr) {
+ std::uninitialized_value_construct(
+ getTrailingObjects<const OpenACCClause *>(),
+ getTrailingObjects<const OpenACCClause *>() + NumClauses);
+ setClauseList(
+ MutableArrayRef(getTrailingObjects<const OpenACCClause *>(), NumClauses));
+}
+
+OpenACCLoopConstruct::OpenACCLoopConstruct(
+ SourceLocation Start, SourceLocation DirLoc, SourceLocation End,
+ ArrayRef<const OpenACCClause *> Clauses, Stmt *Loop)
+ : OpenACCAssociatedStmtConstruct(OpenACCLoopConstructClass,
+ OpenACCDirectiveKind::Loop, Start, DirLoc,
+ End, Loop) {
+ // accept 'nullptr' for the loop. This is diagnosed somewhere, but this gives
+ // us some level of AST fidelity in the error case.
+ assert((Loop == nullptr || isa<ForStmt, CXXForRangeStmt>(Loop)) &&
+ "Associated Loop not a for loop?");
+ // Initialize the trailing storage.
+ std::uninitialized_copy(Clauses.begin(), Clauses.end(),
+ getTrailingObjects<const OpenACCClause *>());
+
+ setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
+ Clauses.size()));
+}
+
+void OpenACCLoopConstruct::setLoop(Stmt *Loop) {
+ assert((isa<ForStmt, CXXForRangeStmt>(Loop)) &&
+ "Associated Loop not a for loop?");
+ setAssociatedStmt(Loop);
+}
+
+OpenACCLoopConstruct *OpenACCLoopConstruct::CreateEmpty(const ASTContext &C,
+ unsigned NumClauses) {
+ void *Mem =
+ C.Allocate(OpenACCLoopConstruct::totalSizeToAlloc<const OpenACCClause *>(
+ NumClauses));
+ auto *Inst = new (Mem) OpenACCLoopConstruct(NumClauses);
+ return Inst;
+}
+
+OpenACCLoopConstruct *
+OpenACCLoopConstruct::Create(const ASTContext &C, SourceLocation BeginLoc,
+ SourceLocation DirLoc, SourceLocation EndLoc,
+ ArrayRef<const OpenACCClause *> Clauses,
+ Stmt *Loop) {
+ void *Mem =
+ C.Allocate(OpenACCLoopConstruct::totalSizeToAlloc<const OpenACCClause *>(
+ Clauses.size()));
+ auto *Inst =
+ new (Mem) OpenACCLoopConstruct(BeginLoc, DirLoc, EndLoc, Clauses, Loop);
return Inst;
}
diff --git a/clang/lib/AST/StmtPrinter.cpp b/clang/lib/AST/StmtPrinter.cpp
index be2d5a2eb6b46..7e030e0551269 100644
--- a/clang/lib/AST/StmtPrinter.cpp
+++ b/clang/lib/AST/StmtPrinter.cpp
@@ -1156,6 +1156,19 @@ void StmtPrinter::VisitOpenACCComputeConstruct(OpenACCComputeConstruct *S) {
PrintStmt(S->getStructuredBlock());
}
+void StmtPrinter::VisitOpenACCLoopConstruct(OpenACCLoopConstruct *S) {
+ Indent() << "#pragma acc loop";
+
+ if (!S->clauses().empty()) {
+ OS << ' ';
+ OpenACCClausePrinter Printer(OS, Policy);
+ Printer.VisitClauseList(S->clauses());
+ }
+ OS << '\n';
+
+ PrintStmt(S->getLoop());
+}
+
//===----------------------------------------------------------------------===//
// Expr printing methods.
//===----------------------------------------------------------------------===//
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index 00b8c43af035c..6d9a76120cfef 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -2605,6 +2605,14 @@ void StmtProfiler::VisitOpenACCComputeConstruct(
P.VisitOpenACCClauseList(S->clauses());
}
+void StmtProfiler::VisitOpenACCLoopConstruct(const OpenACCLoopConstruct *S) {
+ // VisitStmt handles children, so the Loop is handled.
+ VisitStmt(S);
+
+ OpenACCClauseProfiler P{*this};
+ P.VisitOpenACCClauseList(S->clauses());
+}
+
void Stmt::Profile(llvm::FoldingSetNodeID &ID, const ASTContext &Context,
bool Canonical, bool ProfileLambdaExpr) const {
StmtProfilerWithPointers Profiler(ID, Context, Canonical, ProfileLambdaExpr);
diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp
index a0eedc71ea220..194370c1d82e0 100644
--- a/clang/lib/AST/TextNodeDumper.cpp
+++ b/clang/lib/AST/TextNodeDumper.cpp
@@ -2848,3 +2848,10 @@ void TextNodeDumper::VisitHLSLBufferDecl(const HLSLBufferDecl *D) {
void TextNodeDumper::VisitOpenACCConstructStmt(const OpenACCConstructStmt *S) {
OS << " " << S->getDirectiveKind();
}
+void TextNodeDumper::VisitOpenACCLoopConstruct(const OpenACCLoopConstruct *S) {
+
+ if (S->isOrphanedLoopConstruct())
+ OS << " <orphan>";
+ else
+ OS << " parent: " << S->getParentComputeConstruct();
+}
diff --git a/clang/lib/CodeGen/CGStmt.cpp b/clang/lib/CodeGen/CGStmt.cpp
index 99daaa14cf3fe..41ac511c52f51 100644
--- a/clang/lib/CodeGen/CGStmt.cpp
+++ b/clang/lib/CodeGen/CGStmt.cpp
@@ -442,6 +442,9 @@ void CodeGenFunction::EmitStmt(const Stmt *S, ArrayRef<const Attr *> Attrs) {
case Stmt::OpenACCComputeConstructClass:
EmitOpenACCComputeConstruct(cast<OpenACCComputeConstruct>(*S));
break;
+ case Stmt::OpenACCLoopConstructClass:
+ EmitOpenACCLoopConstruct(cast<OpenACCLoopConstruct>(*S));
+ break;
}
}
diff --git a/clang/lib/CodeGen/CodeGenFunction.h b/clang/lib/CodeGen/CodeGenFunction.h
index 45585361a4fc9..5739fbaaa9194 100644
--- a/clang/lib/CodeGen/CodeGenFunction.h
+++ b/clang/lib/CodeGen/CodeGenFunction.h
@@ -4062,6 +4062,13 @@ class CodeGenFunction : public CodeGenTypeCache {
EmitStmt(S.getStructuredBlock());
}
+ void EmitOpenACCLoopConstruct(const OpenACCLoopConstruct &S) {
+ // TODO OpenACC: Implement this. It is currently implemented as a 'no-op',
+ // simply emitting its loop, but in the future we will implement
+ // some sort of IR.
+ EmitStmt(S.getLoop());
+ }
+
//===--------------------------------------------------------------------===//
// LValue Expression Emission
//===--------------------------------------------------------------------===//
diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
index 63afc18783a1f..c7b6763b4dbdd 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -571,6 +571,7 @@ bool doesDirectiveHaveAssociatedStmt(OpenACCDirectiveKind DirKind) {
case OpenACCDirectiveKind::Parallel:
case OpenACCDirectiveKind::Serial:
case OpenACCDirectiveKind::Kernels:
+ case OpenACCDirectiveKind::Loop:
return true;
}
llvm_unreachable("Unhandled directive->assoc stmt");
@@ -1447,13 +1448,14 @@ StmtResult Parser::ParseOpenACCDirectiveStmt() {
return StmtError();
StmtResult AssocStmt;
-
+ SemaOpenACC::AssociatedStmtRAII AssocStmtRAII(getActions().OpenACC(),
+ DirInfo.DirKind);
if (doesDirectiveHaveAssociatedStmt(DirInfo.DirKind)) {
ParsingOpenACCDirectiveRAII DirScope(*this, /*Value=*/false);
ParseScope ACCScope(this, getOpenACCScopeFlags(DirInfo.DirKind));
- AssocStmt = getActions().OpenACC().ActOnAssociatedStmt(DirInfo.DirKind,
- ParseStatement());
+ AssocStmt = getActions().OpenACC().ActOnAssociatedStmt(
+ DirInfo.StartLoc, DirInfo.DirKind, ParseStatement());
}
return getActions().OpenACC().ActOnEndStmtDirective(
diff --git a/clang/lib/Sema/SemaExceptionSpec.cpp b/clang/lib/Sema/SemaExceptionSpec.cpp
index 41bf273d12f2f..17acfca6b0112 100644
--- a/clang/lib/Sema/SemaExceptionSpec.cpp
+++ b/clang/lib/Sema/SemaExceptionSpec.cpp
@@ -1425,6 +1425,7 @@ CanThrowResult Sema::canThrow(const Stmt *S) {
// Most statements can throw if any substatement can throw.
case Stmt::OpenACCComputeConstructClass:
+ case Stmt::OpenACCLoopConstructClass:
case Stmt::AttributedStmtClass:
case Stmt::BreakStmtClass:
case Stmt::CapturedStmtClass:
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 15239f4f35c39..f21bc945b451f 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -33,6 +33,7 @@ bool diagnoseConstructAppertainment(SemaOpenACC &S, OpenACCDirectiveKind K,
case OpenACCDirectiveKind::Parallel:
case OpenACCDirectiveKind::Serial:
case OpenACCDirectiveKind::Kernels:
+ case OpenACCDirectiveKind::Loop:
if (!IsStmt)
return S.Diag(StartLoc, diag::err_acc_construct_appertainment) << K;
break;
@@ -298,6 +299,30 @@ bool checkValidAfterDeviceType(
SemaOpenACC::SemaOpenACC(Sema &S) : SemaBase(S) {}
+SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(SemaOpenACC &S,
+ OpenACCDirectiveKind DK)
+ : SemaRef(S), WasInsideComputeConstruct(S.InsideComputeConstruct),
+ DirKind(DK) {
+ // Compute constructs end up taking their 'loop'.
+ if (DirKind == OpenACCDirectiveKind::Parallel ||
+ DirKind == OpenACCDirectiveKind::Serial ||
+ DirKind == OpenACCDirectiveKind::Kernels) {
+ SemaRef.InsideComputeConstruct = true;
+ SemaRef.ParentlessLoopConstructs.swap(ParentlessLoopConstructs);
+ }
+}
+
+SemaOpenACC::AssociatedStmtRAII::~AssociatedStmtRAII() {
+ SemaRef.InsideComputeConstruct = WasInsideComputeConstruct;
+ if (DirKind == OpenACCDirectiveKind::Parallel ||
+ DirKind == OpenACCDirectiveKind::Serial ||
+ DirKind == OpenACCDirectiveKind::Kernels) {
+ assert(SemaRef.ParentlessLoopConstructs.empty() &&
+ "Didn't consume loop construct list?");
+ SemaRef.ParentlessLoopConstructs.swap(ParentlessLoopConstructs);
+ }
+}
+
OpenACCClause *
SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses,
OpenACCParsedClause &Clause) {
@@ -855,6 +880,7 @@ void SemaOpenACC::ActOnConstruct(OpenACCDirectiveKind K,
case OpenACCDirectiveKind::Parallel:
case OpenACCDirectiveKind::Serial:
case OpenACCDirectiveKind::Kernels:
+ case OpenACCDirectiveKind::Loop:
// Nothing to do here, there is no real legalization that needs to happen
// here as these constructs do not take any arguments.
break;
@@ -1276,16 +1302,34 @@ StmtResult SemaOpenACC::ActOnEndStmtDirective(OpenACCDirectiveKind K,
return StmtError();
case OpenACCDirectiveKind::Parallel:
case OpenACCDirectiveKind::Serial:
- case OpenACCDirectiveKind::Kernels:
- // TODO OpenACC: Add clauses to the construct here.
- return OpenACCComputeConstruct::Create(
+ case OpenACCDirectiveKind::Kernels: {
+ auto *ComputeConstruct = OpenACCComputeConstruct::Create(
getASTContext(), K, StartLoc, DirLoc, EndLoc, Clauses,
+ AssocStmt.isUsable() ? AssocStmt.get() : nullptr,
+ ParentlessLoopConstructs);
+
+ ParentlessLoopConstructs.clear();
+ return ComputeConstruct;
+ }
+ case OpenACCDirectiveKind::Loop: {
+ auto *LoopConstruct = OpenACCLoopConstruct::Create(
+ getASTContext(), StartLoc, DirLoc, EndLoc, Clauses,
AssocStmt.isUsable() ? AssocStmt.get() : nullptr);
+
+ // If we are in the scope of a compute construct, add this to the list of
+ // loop constructs that need assigning to the next closing compute
+ // construct.
+ if (InsideComputeConstruct)
+ ParentlessLoopConstructs.push_back(LoopConstruct);
+
+ return LoopConstruct;
+ }
}
llvm_unreachable("Unhandled case in directive handling?");
}
-StmtResult SemaOpenACC::ActOnAssociatedStmt(OpenACCDirectiveKind K,
+StmtResult SemaOpenACC::ActOnAssociatedStmt(SourceLocation DirectiveLoc,
+ OpenACCDirectiveKind K,
StmtResult AssocStmt) {
switch (K) {
default:
@@ -1303,6 +1347,14 @@ StmtResult SemaOpenACC::ActOnAssociatedStmt(OpenACCDirectiveKind K,
// an interpretation of it is to allow this and treat the initializer as
// the 'structured block'.
return AssocStmt;
+ case OpenACCDirectiveKind::Loop:
+ if (AssocStmt.isUsable() &&
+ !isa<CXXForRangeStmt, ForStmt>(AssocStmt.get())) {
+ Diag(AssocStmt.get()->getBeginLoc(), diag::err_acc_loop_not_for_loop);
+ Diag(DirectiveLoc, diag::note_acc_construct_here) << K;
+ return StmtError();
+ }
+ return AssocStmt;
}
llvm_unreachable("Invalid associated statement application");
}
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index efba99b85b0fb..ab6a59e61d370 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -4041,6 +4041,15 @@ class TreeTransform {
EndLoc, Clauses, StrBlock);
}
+ StmtResult RebuildOpenACCLoopConstruct(SourceLocation BeginLoc,
+ SourceLocation DirLoc,
+ SourceLocation EndLoc,
+ ArrayRef<OpenACCClause *> Clauses,
+ StmtResult Loop) {
+ return getSema().OpenACC().ActOnEndStmtDirective(
+ OpenACCDirectiveKind::Loop, BeginLoc, DirLoc, EndLoc, Clauses, Loop);
+ }
+
private:
TypeLoc TransformTypeInObjectScope(TypeLoc TL,
QualType ObjectType,
@@ -11541,8 +11550,6 @@ template <typename Derived>
StmtResult TreeTransform<Derived>::TransformOpenACCComputeConstruct(
OpenACCComputeConstruct *C) {
getSema().OpenACC().ActOnConstruct(C->getDirectiveKind(), C->getBeginLoc());
- // FIXME: When implementing this for constructs that can take arguments, we
- // should do Sema for them here.
if (getSema().OpenACC().ActOnStartStmtDirective(C->getDirectiveKind(),
C->getBeginLoc()))
@@ -11551,17 +11558,44 @@ StmtResult TreeTransform<Derived>::TransformOpenACCComputeConstruct(
llvm::SmallVector<OpenACCClause *> TransformedClauses =
getDerived().TransformOpenACCClauseList(C->getDirectiveKind(),
C->clauses());
-
// Transform Structured Block.
+ SemaOpenACC::AssociatedStmtRAII AssocStmtRAII(getSema().OpenACC(),
+ C->getDirectiveKind());
StmtResult StrBlock = getDerived().TransformStmt(C->getStructuredBlock());
- StrBlock =
- getSema().OpenACC().ActOnAssociatedStmt(C->getDirectiveKind(), StrBlock);
+ StrBlock = getSema().OpenACC().ActOnAssociatedStmt(
+ C->getBeginLoc(), C->getDirectiveKind(), StrBlock);
return getDerived().RebuildOpenACCComputeConstruct(
C->getDirectiveKind(), C->getBeginLoc(), C->getDirectiveLoc(),
C->getEndLoc(), TransformedClauses, StrBlock);
}
+template <typename Derived>
+StmtResult
+TreeTransform<Derived>::TransformOpenACCLoopConstruct(OpenACCLoopConstruct *C) {
+
+ getSema().OpenACC().ActOnConstruct(C->getDirectiveKind(), C->getBeginLoc());
+
+ if (getSema().OpenACC().ActOnStartStmtDirective(C->getDirectiveKind(),
+ C->getBeginLoc()))
+ return StmtError();
+
+ llvm::SmallVector<OpenACCClause *> TransformedClauses =
+ getDerived().TransformOpenACCClauseList(C->getDirectiveKind(),
+ C->clauses());
+
+ // Transform Loop.
+ SemaOpenACC::AssociatedStmtRAII AssocStmtRAII(getSema().OpenACC(),
+ C->getDirectiveKind());
+ StmtResult Loop = getDerived().TransformStmt(C->getLoop());
+ Loop = getSema().OpenACC().ActOnAssociatedStmt(C->getBeginLoc(),
+ C->getDirectiveKind(), Loop);
+
+ return getDerived().RebuildOpenACCLoopConstruct(
+ C->getBeginLoc(), C->getDirectiveLoc(), C->getEndLoc(),
+ TransformedClauses, Loop);
+}
+
//===----------------------------------------------------------------------===//
// Expression transformation
//===----------------------------------------------------------------------===//
diff --git a/clang/lib/Serialization/ASTReaderStmt.cpp b/clang/lib/Serialization/ASTReaderStmt.cpp
index bea2b94989107..67ef170251914 100644
--- a/clang/lib/Serialization/ASTReaderStmt.cpp
+++ b/clang/lib/Serialization/ASTReaderStmt.cpp
@@ -2810,6 +2810,12 @@ void ASTStmtReader::VisitOpenACCAssociatedStmtConstruct(
void ASTStmtReader::VisitOpenACCComputeConstruct(OpenACCComputeConstruct *S) {
VisitStmt(S);
VisitOpenACCAssociatedStmtConstruct(S);
+ S->findAndSetChildLoops();
+}
+
+void ASTStmtReader::VisitOpenACCLoopConstruct(OpenACCLoopConstruct *S) {
+ VisitStmt(S);
+ VisitOpenACCAssociatedStmtConstruct(S);
}
//===----------------------------------------------------------------------===//
@@ -4235,6 +4241,11 @@ Stmt *ASTReader::ReadStmtFromStream(ModuleFile &F) {
S = OpenACCComputeConstruct::CreateEmpty(Context, NumClauses);
break;
}
+ case STMT_OPENACC_LOOP_CONSTRUCT: {
+ unsigned NumClauses = Record[ASTStmtReader::NumStmtFields];
+ S = OpenACCLoopConstruct::CreateEmpty(Context, NumClauses);
+ break;
+ }
case EXPR_REQUIRES:
unsigned numLocalParameters = Record[ASTStmtReader::NumExprFields];
unsigned numRequirement = Record[ASTStmtReader::NumExprFields + 1];
diff --git a/clang/lib/Serialization/ASTWriterStmt.cpp b/clang/lib/Serialization/ASTWriterStmt.cpp
index 3c586b270fbf4..1a98e30e0f89f 100644
--- a/clang/lib/Serialization/ASTWriterStmt.cpp
+++ b/clang/lib/Serialization/ASTWriterStmt.cpp
@@ -2863,6 +2863,12 @@ void ASTStmtWriter::VisitOpenACCComputeConstruct(OpenACCComputeConstruct *S) {
Code = serialization::STMT_OPENACC_COMPUTE_CONSTRUCT;
}
+void ASTStmtWriter::VisitOpenACCLoopConstruct(OpenACCLoopConstruct *S) {
+ VisitStmt(S);
+ VisitOpenACCAssociatedStmtConstruct(S);
+ Code = serialization::STMT_OPENACC_LOOP_CONSTRUCT;
+}
+
//===----------------------------------------------------------------------===//
// ASTWriter Implementation
//===----------------------------------------------------------------------===//
diff --git a/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp b/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
index 793f3a63ea29e..290d96611d466 100644
--- a/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
+++ b/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
@@ -1822,6 +1822,7 @@ void ExprEngine::Visit(const Stmt *S, ExplodedNode *Pred,
case Stmt::OMPTargetParallelGenericLoopDirectiveClass:
case Stmt::CapturedStmtClass:
case Stmt::OpenACCComputeConstructClass:
+ case Stmt::OpenACCLoopConstructClass:
case Stmt::OMPUnrollDirectiveClass:
case Stmt::OMPMetaDirectiveClass: {
const ExplodedNode *node = Bldr.generateSink(S, Pred, Pred->getState());
diff --git a/clang/test/AST/ast-print-openacc-loop-construct.cpp b/clang/test/AST/ast-print-openacc-loop-construct.cpp
new file mode 100644
index 0000000000000..21c92b17317ef
--- /dev/null
+++ b/clang/test/AST/ast-print-openacc-loop-construct.cpp
@@ -0,0 +1,9 @@
+// RUN: %clang_cc1 -fopenacc -Wno-openacc-deprecated-clause-alias -ast-print %s -o - | FileCheck %s
+
+void foo() {
+// CHECK: #pragma acc loop
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc loop
+ for(;;);
+}
diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c
index 49e749feb2ec7..cb118f69fb447 100644
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@@ -37,23 +37,23 @@ void func() {
// expected-warning at +1{{OpenACC construct 'host_data' not yet implemented, pragma ignored}}
#pragma acc host_data if_present, if_present
- // expected-warning at +4{{OpenACC clause 'seq' not yet implemented, clause ignored}}
- // expected-warning at +3{{OpenACC clause 'independent' not yet implemented, clause ignored}}
- // expected-warning at +2{{OpenACC clause 'auto' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-warning at +3{{OpenACC clause 'seq' not yet implemented, clause ignored}}
+ // expected-warning at +2{{OpenACC clause 'independent' not yet implemented, clause ignored}}
+ // expected-warning at +1{{OpenACC clause 'auto' not yet implemented, clause ignored}}
#pragma acc loop seq independent auto
+ for(;;){}
- // expected-warning at +4{{OpenACC clause 'seq' not yet implemented, clause ignored}}
- // expected-warning at +3{{OpenACC clause 'independent' not yet implemented, clause ignored}}
- // expected-warning at +2{{OpenACC clause 'auto' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-warning at +3{{OpenACC clause 'seq' not yet implemented, clause ignored}}
+ // expected-warning at +2{{OpenACC clause 'independent' not yet implemented, clause ignored}}
+ // expected-warning at +1{{OpenACC clause 'auto' not yet implemented, clause ignored}}
#pragma acc loop seq, independent auto
+ for(;;){}
- // expected-warning at +4{{OpenACC clause 'seq' not yet implemented, clause ignored}}
- // expected-warning at +3{{OpenACC clause 'independent' not yet implemented, clause ignored}}
- // expected-warning at +2{{OpenACC clause 'auto' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-warning at +3{{OpenACC clause 'seq' not yet implemented, clause ignored}}
+ // expected-warning at +2{{OpenACC clause 'independent' not yet implemented, clause ignored}}
+ // expected-warning at +1{{OpenACC clause 'auto' not yet implemented, clause ignored}}
#pragma acc loop seq independent, auto
+ for(;;){}
// expected-warning at +4{{OpenACC clause 'seq' not yet implemented, clause ignored}}
// expected-warning at +3{{OpenACC clause 'independent' not yet implemented, clause ignored}}
@@ -67,65 +67,57 @@ void func() {
// expected-warning at +2{{OpenACC clause 'auto' not yet implemented, clause ignored}}
// expected-warning at +1{{OpenACC construct 'serial loop' not yet implemented, pragma ignored}}
#pragma acc serial loop seq, independent auto
- {}
+ for(;;){}
// expected-warning at +4{{OpenACC clause 'seq' not yet implemented, clause ignored}}
// expected-warning at +3{{OpenACC clause 'independent' not yet implemented, clause ignored}}
// expected-warning at +2{{OpenACC clause 'auto' not yet implemented, clause ignored}}
// expected-warning at +1{{OpenACC construct 'parallel loop' not yet implemented, pragma ignored}}
#pragma acc parallel loop seq independent, auto
- {}
+ for(;;){}
- // expected-error at +2{{expected identifier}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +1{{expected identifier}}
#pragma acc loop , seq
+ for(;;){}
- // expected-error at +3{{expected identifier}}
- // expected-warning at +2{{OpenACC clause 'seq' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +2{{expected identifier}}
+ // expected-warning at +1{{OpenACC clause 'seq' not yet implemented, clause ignored}}
#pragma acc loop seq,
+ for(;;){}
- // expected-error at +2{{expected '('}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +1{{expected '('}}
#pragma acc loop collapse
for(;;){}
- // expected-error at +2{{expected expression}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +1{{expected expression}}
#pragma acc loop collapse()
for(;;){}
- // expected-error at +3{{invalid tag 'unknown' on 'collapse' clause}}
- // expected-error at +2{{expected expression}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +2{{invalid tag 'unknown' on 'collapse' clause}}
+ // expected-error at +1{{expected expression}}
#pragma acc loop collapse(unknown:)
for(;;){}
- // expected-error at +2{{expected expression}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +1{{expected expression}}
#pragma acc loop collapse(force:)
for(;;){}
- // expected-error at +3{{invalid tag 'unknown' on 'collapse' clause}}
- // expected-warning at +2{{OpenACC clause 'collapse' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +2{{invalid tag 'unknown' on 'collapse' clause}}
+ // expected-warning at +1{{OpenACC clause 'collapse' not yet implemented, clause ignored}}
#pragma acc loop collapse(unknown:5)
for(;;){}
- // expected-warning at +2{{OpenACC clause 'collapse' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-warning at +1{{OpenACC clause 'collapse' not yet implemented, clause ignored}}
#pragma acc loop collapse(force:5)
for(;;){}
- // expected-warning at +2{{OpenACC clause 'collapse' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-warning at +1{{OpenACC clause 'collapse' not yet implemented, clause ignored}}
#pragma acc loop collapse(5)
for(;;){}
- // expected-error at +3{{expected ')'}}
- // expected-note at +2{{to match this '('}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +2{{expected ')'}}
+ // expected-note at +1{{to match this '('}}
#pragma acc loop collapse(5, 6)
for(;;){}
}
@@ -989,108 +981,108 @@ void IntExprParsing() {
#pragma acc set default_async(returns_int())
- // expected-warning at +2{{OpenACC clause 'vector' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-warning at +1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
#pragma acc loop vector
- // expected-error at +2{{expected expression}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ for(;;);
+ // expected-error at +1{{expected expression}}
#pragma acc loop vector()
- // expected-error at +3{{invalid tag 'invalid' on 'vector' clause}}
- // expected-error at +2{{expected expression}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ for(;;);
+ // expected-error at +2{{invalid tag 'invalid' on 'vector' clause}}
+ // expected-error at +1{{expected expression}}
#pragma acc loop vector(invalid:)
- // expected-error at +3{{invalid tag 'invalid' on 'vector' clause}}
- // expected-warning at +2{{OpenACC clause 'vector' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ for(;;);
+ // expected-error at +2{{invalid tag 'invalid' on 'vector' clause}}
+ // expected-warning at +1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
#pragma acc loop vector(invalid:5)
- // expected-error at +2{{expected expression}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ for(;;);
+ // expected-error at +1{{expected expression}}
#pragma acc loop vector(length:)
- // expected-error at +3{{invalid tag 'num' on 'vector' clause}}
- // expected-error at +2{{expected expression}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ for(;;);
+ // expected-error at +2{{invalid tag 'num' on 'vector' clause}}
+ // expected-error at +1{{expected expression}}
#pragma acc loop vector(num:)
- // expected-error at +3{{expected ')'}}
- // expected-note at +2{{to match this '('}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ for(;;);
+ // expected-error at +2{{expected ')'}}
+ // expected-note at +1{{to match this '('}}
#pragma acc loop vector(5, 4)
- // expected-error at +3{{expected ')'}}
- // expected-note at +2{{to match this '('}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ for(;;);
+ // expected-error at +2{{expected ')'}}
+ // expected-note at +1{{to match this '('}}
#pragma acc loop vector(length:6,4)
- // expected-error at +4{{invalid tag 'num' on 'vector' clause}}
- // expected-error at +3{{expected ')'}}
- // expected-note at +2{{to match this '('}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ for(;;);
+ // expected-error at +3{{invalid tag 'num' on 'vector' clause}}
+ // expected-error at +2{{expected ')'}}
+ // expected-note at +1{{to match this '('}}
#pragma acc loop vector(num:6,4)
- // expected-warning at +2{{OpenACC clause 'vector' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ for(;;);
+ // expected-warning at +1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
#pragma acc loop vector(5)
- // expected-error at +3{{invalid tag 'num' on 'vector' clause}}
- // expected-warning at +2{{OpenACC clause 'vector' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ for(;;);
+ // expected-error at +2{{invalid tag 'num' on 'vector' clause}}
+ // expected-warning at +1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
#pragma acc loop vector(num:5)
- // expected-warning at +2{{OpenACC clause 'vector' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ for(;;);
+ // expected-warning at +1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
#pragma acc loop vector(length:5)
- // expected-warning at +2{{OpenACC clause 'vector' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ for(;;);
+ // expected-warning at +1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
#pragma acc loop vector(returns_int())
- // expected-warning at +2{{OpenACC clause 'vector' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ for(;;);
+ // expected-warning at +1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
#pragma acc loop vector(length:returns_int())
+ for(;;);
- // expected-warning at +2{{OpenACC clause 'worker' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-warning at +1{{OpenACC clause 'worker' not yet implemented, clause ignored}}
#pragma acc loop worker
- // expected-error at +2{{expected expression}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ for(;;);
+ // expected-error at +1{{expected expression}}
#pragma acc loop worker()
- // expected-error at +3{{invalid tag 'invalid' on 'worker' clause}}
- // expected-error at +2{{expected expression}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ for(;;);
+ // expected-error at +2{{invalid tag 'invalid' on 'worker' clause}}
+ // expected-error at +1{{expected expression}}
#pragma acc loop worker(invalid:)
- // expected-error at +3{{invalid tag 'invalid' on 'worker' clause}}
- // expected-warning at +2{{OpenACC clause 'worker' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ for(;;);
+ // expected-error at +2{{invalid tag 'invalid' on 'worker' clause}}
+ // expected-warning at +1{{OpenACC clause 'worker' not yet implemented, clause ignored}}
#pragma acc loop worker(invalid:5)
- // expected-error at +2{{expected expression}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ for(;;);
+ // expected-error at +1{{expected expression}}
#pragma acc loop worker(num:)
- // expected-error at +3{{invalid tag 'length' on 'worker' clause}}
- // expected-error at +2{{expected expression}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ for(;;);
+ // expected-error at +2{{invalid tag 'length' on 'worker' clause}}
+ // expected-error at +1{{expected expression}}
#pragma acc loop worker(length:)
- // expected-error at +3{{expected ')'}}
- // expected-note at +2{{to match this '('}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ for(;;);
+ // expected-error at +2{{expected ')'}}
+ // expected-note at +1{{to match this '('}}
#pragma acc loop worker(5, 4)
- // expected-error at +3{{expected ')'}}
- // expected-note at +2{{to match this '('}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ for(;;);
+ // expected-error at +2{{expected ')'}}
+ // expected-note at +1{{to match this '('}}
#pragma acc loop worker(num:6,4)
- // expected-error at +4{{invalid tag 'length' on 'worker' clause}}
- // expected-error at +3{{expected ')'}}
- // expected-note at +2{{to match this '('}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ for(;;);
+ // expected-error at +3{{invalid tag 'length' on 'worker' clause}}
+ // expected-error at +2{{expected ')'}}
+ // expected-note at +1{{to match this '('}}
#pragma acc loop worker(length:6,4)
- // expected-warning at +2{{OpenACC clause 'worker' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ for(;;);
+ // expected-warning at +1{{OpenACC clause 'worker' not yet implemented, clause ignored}}
#pragma acc loop worker(5)
- // expected-error at +3{{invalid tag 'length' on 'worker' clause}}
- // expected-warning at +2{{OpenACC clause 'worker' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ for(;;);
+ // expected-error at +2{{invalid tag 'length' on 'worker' clause}}
+ // expected-warning at +1{{OpenACC clause 'worker' not yet implemented, clause ignored}}
#pragma acc loop worker(length:5)
- // expected-warning at +2{{OpenACC clause 'worker' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ for(;;);
+ // expected-warning at +1{{OpenACC clause 'worker' not yet implemented, clause ignored}}
#pragma acc loop worker(num:5)
- // expected-warning at +2{{OpenACC clause 'worker' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ for(;;);
+ // expected-warning at +1{{OpenACC clause 'worker' not yet implemented, clause ignored}}
#pragma acc loop worker(returns_int())
- // expected-error at +3{{invalid tag 'length' on 'worker' clause}}
- // expected-warning at +2{{OpenACC clause 'worker' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ for(;;);
+ // expected-error at +2{{invalid tag 'length' on 'worker' clause}}
+ // expected-warning at +1{{OpenACC clause 'worker' not yet implemented, clause ignored}}
#pragma acc loop worker(length:returns_int())
+ for(;;);
}
void device_type() {
@@ -1236,238 +1228,196 @@ void AsyncArgument() {
void Tile() {
int* Foo;
- // expected-error at +2{{expected '('}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +1{{expected '('}}
#pragma acc loop tile
for(;;){}
- // expected-error at +4{{expected expression}}
- // expected-error at +3{{expected ')'}}
- // expected-note at +2{{to match this '('}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +3{{expected expression}}
+ // expected-error at +2{{expected ')'}}
+ // expected-note at +1{{to match this '('}}
#pragma acc loop tile(
for(;;){}
- // expected-error at +3{{expected expression}}
- // expected-warning at +2{{OpenACC clause 'tile' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +2{{expected expression}}
+ // expected-warning at +1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
#pragma acc loop tile()
for(;;){}
- // expected-error at +4{{expected expression}}
- // expected-error at +3{{expected ')'}}
- // expected-note at +2{{to match this '('}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +3{{expected expression}}
+ // expected-error at +2{{expected ')'}}
+ // expected-note at +1{{to match this '('}}
#pragma acc loop tile(,
for(;;){}
- // expected-error at +3{{expected expression}}
- // expected-warning at +2{{OpenACC clause 'tile' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +2{{expected expression}}
+ // expected-warning at +1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
#pragma acc loop tile(,)
for(;;){}
- // expected-error at +3{{use of undeclared identifier 'invalid'}}
- // expected-warning at +2{{OpenACC clause 'tile' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +2{{use of undeclared identifier 'invalid'}}
+ // expected-warning at +1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
#pragma acc loop tile(returns_int(), *, invalid, *)
for(;;){}
- // expected-error at +3{{expected expression}}
- // expected-warning at +2{{OpenACC clause 'tile' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +2{{expected expression}}
+ // expected-warning at +1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
#pragma acc loop tile(returns_int() *, Foo, *)
for(;;){}
- // expected-error at +3{{indirection requires pointer operand ('int' invalid)}}
- // expected-warning at +2{{OpenACC clause 'tile' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +2{{indirection requires pointer operand ('int' invalid)}}
+ // expected-warning at +1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
#pragma acc loop tile(* returns_int() , *)
for(;;){}
- // expected-warning at +2{{OpenACC clause 'tile' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-warning at +1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
#pragma acc loop tile(*)
for(;;){}
- // expected-warning at +2{{OpenACC clause 'tile' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-warning at +1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
#pragma acc loop tile(*Foo, *Foo)
for(;;){}
- // expected-warning at +2{{OpenACC clause 'tile' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-warning at +1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
#pragma acc loop tile(5)
for(;;){}
- // expected-warning at +2{{OpenACC clause 'tile' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-warning at +1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
#pragma acc loop tile(*, 5)
for(;;){}
- // expected-warning at +2{{OpenACC clause 'tile' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-warning at +1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
#pragma acc loop tile(5, *)
for(;;){}
- // expected-warning at +2{{OpenACC clause 'tile' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-warning at +1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
#pragma acc loop tile(5, *, 3, *)
for(;;){}
}
void Gang() {
- // expected-warning at +2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
#pragma acc loop gang
for(;;){}
- // expected-error at +4{{expected expression}}
- // expected-error at +3{{expected ')'}}
- // expected-note at +2{{to match this '('}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +3{{expected expression}}
+ // expected-error at +2{{expected ')'}}
+ // expected-note at +1{{to match this '('}}
#pragma acc loop gang(
for(;;){}
- // expected-error at +3{{expected expression}}
- // expected-warning at +2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +2{{expected expression}}
+ // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
#pragma acc loop gang()
for(;;){}
- // expected-error at +3{{expected expression}}
- // expected-warning at +2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +2{{expected expression}}
+ // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
#pragma acc loop gang(5, *)
for(;;){}
- // expected-error at +3{{expected expression}}
- // expected-warning at +2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +2{{expected expression}}
+ // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
#pragma acc loop gang(*)
for(;;){}
- // expected-error at +3{{expected expression}}
- // expected-warning at +2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +2{{expected expression}}
+ // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
#pragma acc loop gang(5, num:*)
for(;;){}
- // expected-error at +3{{expected expression}}
- // expected-warning at +2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +2{{expected expression}}
+ // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
#pragma acc loop gang(num:5, *)
for(;;){}
- // expected-error at +3{{expected expression}}
- // expected-warning at +2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +2{{expected expression}}
+ // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
#pragma acc loop gang(num:5, num:*)
for(;;){}
- // expected-error at +3{{expected expression}}
- // expected-warning at +2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +2{{expected expression}}
+ // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
#pragma acc loop gang(num:*)
for(;;){}
- // expected-warning at +2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
#pragma acc loop gang(dim:5)
for(;;){}
- // expected-error at +3{{expected expression}}
- // expected-warning at +2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +2{{expected expression}}
+ // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
#pragma acc loop gang(dim:5, dim:*)
for(;;){}
- // expected-error at +3{{expected expression}}
- // expected-warning at +2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +2{{expected expression}}
+ // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
#pragma acc loop gang(dim:*)
for(;;){}
- // expected-warning at +2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
#pragma acc loop gang(static:*)
for(;;){}
- // expected-warning at +2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
#pragma acc loop gang(static:*, static:5)
for(;;){}
- // expected-warning at +2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
#pragma acc loop gang(static:*, 5)
for(;;){}
- // expected-warning at +2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
#pragma acc loop gang(static:45, 5)
for(;;){}
- // expected-error at +4{{expected expression}}
- // expected-error at +3{{expected ')'}}
- // expected-note at +2{{to match this '('}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +3{{expected expression}}
+ // expected-error at +2{{expected ')'}}
+ // expected-note at +1{{to match this '('}}
#pragma acc loop gang(static:45,
for(;;){}
- // expected-error at +3{{expected ')'}}
- // expected-note at +2{{to match this '('}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +2{{expected ')'}}
+ // expected-note at +1{{to match this '('}}
#pragma acc loop gang(static:45
for(;;){}
- // expected-error at +4{{expected expression}}
- // expected-error at +3{{expected ')'}}
- // expected-note at +2{{to match this '('}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +3{{expected expression}}
+ // expected-error at +2{{expected ')'}}
+ // expected-note at +1{{to match this '('}}
#pragma acc loop gang(static:*,
for(;;){}
- // expected-error at +3{{expected ')'}}
- // expected-note at +2{{to match this '('}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +2{{expected ')'}}
+ // expected-note at +1{{to match this '('}}
#pragma acc loop gang(static:*
for(;;){}
- // expected-error at +4{{expected expression}}
- // expected-error at +3{{expected ')'}}
- // expected-note at +2{{to match this '('}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +3{{expected expression}}
+ // expected-error at +2{{expected ')'}}
+ // expected-note at +1{{to match this '('}}
#pragma acc loop gang(45,
for(;;){}
- // expected-error at +3{{expected ')'}}
- // expected-note at +2{{to match this '('}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +2{{expected ')'}}
+ // expected-note at +1{{to match this '('}}
#pragma acc loop gang(45
for(;;){}
- // expected-error at +4{{expected expression}}
- // expected-error at +3{{expected ')'}}
- // expected-note at +2{{to match this '('}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +3{{expected expression}}
+ // expected-error at +2{{expected ')'}}
+ // expected-note at +1{{to match this '('}}
#pragma acc loop gang(num:45,
for(;;){}
- // expected-error at +3{{expected ')'}}
- // expected-note at +2{{to match this '('}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +2{{expected ')'}}
+ // expected-note at +1{{to match this '('}}
#pragma acc loop gang(num:45
for(;;){}
- // expected-error at +4{{expected expression}}
- // expected-error at +3{{expected ')'}}
- // expected-note at +2{{to match this '('}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +3{{expected expression}}
+ // expected-error at +2{{expected ')'}}
+ // expected-note at +1{{to match this '('}}
#pragma acc loop gang(dim:45,
for(;;){}
- // expected-error at +3{{expected ')'}}
- // expected-note at +2{{to match this '('}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +2{{expected ')'}}
+ // expected-note at +1{{to match this '('}}
#pragma acc loop gang(dim:45
for(;;){}
- // expected-warning at +2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
#pragma acc loop gang(static:*, dim:returns_int(), 5)
for(;;){}
- // expected-warning at +2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
#pragma acc loop gang(num: 32, static:*, dim:returns_int(), 5)
for(;;){}
diff --git a/clang/test/ParserOpenACC/parse-clauses.cpp b/clang/test/ParserOpenACC/parse-clauses.cpp
index 702eb75ca8902..b7e252e892bea 100644
--- a/clang/test/ParserOpenACC/parse-clauses.cpp
+++ b/clang/test/ParserOpenACC/parse-clauses.cpp
@@ -2,13 +2,11 @@
template<unsigned I, typename T>
void templ() {
- // expected-warning at +2{{OpenACC clause 'collapse' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-warning at +1{{OpenACC clause 'collapse' not yet implemented, clause ignored}}
#pragma acc loop collapse(I)
for(;;){}
- // expected-warning at +2{{OpenACC clause 'collapse' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-warning at +1{{OpenACC clause 'collapse' not yet implemented, clause ignored}}
#pragma acc loop collapse(T::value)
for(;;){}
diff --git a/clang/test/ParserOpenACC/parse-constructs.c b/clang/test/ParserOpenACC/parse-constructs.c
index ecedfd9e9e6d6..ea75360cc1351 100644
--- a/clang/test/ParserOpenACC/parse-constructs.c
+++ b/clang/test/ParserOpenACC/parse-constructs.c
@@ -82,8 +82,7 @@ void func() {
// expected-warning at +1{{OpenACC construct 'host_data' not yet implemented, pragma ignored}}
#pragma acc host_data clause list
for(;;){}
- // expected-error at +2{{invalid OpenACC clause 'clause'}}
- // expected-warning at +1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+ // expected-error at +1{{invalid OpenACC clause 'clause'}}
#pragma acc loop clause list
for(;;){}
// expected-error at +1{{invalid OpenACC clause 'invalid'}}
diff --git a/clang/test/SemaOpenACC/compute-construct-default-clause.c b/clang/test/SemaOpenACC/compute-construct-default-clause.c
index b1235fcca1f6a..df61357652551 100644
--- a/clang/test/SemaOpenACC/compute-construct-default-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-default-clause.c
@@ -43,10 +43,9 @@ void SingleOnly() {
#pragma acc data default(none)
while(0);
- // expected-warning at +2{{OpenACC construct 'loop' not yet implemented}}
// expected-error at +1{{OpenACC 'default' clause is not valid on 'loop' directive}}
#pragma acc loop default(none)
- while(0);
+ for(;;);
// expected-warning at +2{{OpenACC construct 'wait' not yet implemented}}
// expected-error at +1{{OpenACC 'default' clause is not valid on 'wait' directive}}
diff --git a/clang/test/SemaOpenACC/loop-ast.cpp b/clang/test/SemaOpenACC/loop-ast.cpp
new file mode 100644
index 0000000000000..a55ca0b6a804c
--- /dev/null
+++ b/clang/test/SemaOpenACC/loop-ast.cpp
@@ -0,0 +1,164 @@
+
+// RUN: %clang_cc1 %s -fopenacc -ast-dump | FileCheck %s
+
+// Test this with PCH.
+// RUN: %clang_cc1 %s -fopenacc -emit-pch -o %t %s
+// RUN: %clang_cc1 %s -fopenacc -include-pch %t -ast-dump-all | FileCheck %s
+
+#ifndef PCH_HELPER
+#define PCH_HELPER
+
+void NormalFunc() {
+ // CHECK-LABEL: NormalFunc
+ // CHECK-NEXT: CompoundStmt
+
+#pragma acc loop
+ for(;;);
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} <orphan>
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+
+ int array[5];
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl
+#pragma acc loop
+ for(auto x : array){}
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} <orphan>
+ // CHECK-NEXT: CXXForRangeStmt
+ // CHECK: CompoundStmt
+
+#pragma acc parallel
+ // CHECK-NEXT: OpenACCComputeConstruct {{.*}}parallel
+ // CHECK-NEXT: CompoundStmt
+ {
+#pragma acc parallel
+ // CHECK-NEXT: OpenACCComputeConstruct [[PAR_ADDR:[0-9a-fx]+]] {{.*}}parallel
+ // CHECK-NEXT: CompoundStmt
+ {
+#pragma acc loop
+ for(;;);
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: [[PAR_ADDR]]
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+ }
+ }
+}
+
+template<typename T>
+void TemplFunc() {
+ // CHECK-LABEL: FunctionTemplateDecl {{.*}}TemplFunc
+ // CHECK-NEXT: TemplateTypeParmDecl
+ // CHECK-NEXT: FunctionDecl{{.*}}TemplFunc
+ // CHECK-NEXT: CompoundStmt
+
+#pragma acc loop
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} <orphan>
+ for(typename T::type t = 0; t < 5;++t) {
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} referenced t 'typename T::type'
+ // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 0
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: BinaryOperator{{.*}} '<dependent type>' '<'
+ // CHECK-NEXT: DeclRefExpr {{.*}} 'typename T::type' lvalue Var
+ // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 5
+ // CHECK-NEXT: UnaryOperator{{.*}} '<dependent type>' lvalue prefix '++'
+ // CHECK-NEXT: DeclRefExpr {{.*}} 'typename T::type' lvalue Var
+ // CHECK-NEXT: CompoundStmt
+ typename T::type I;
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} I 'typename T::type'
+
+ }
+
+#pragma acc parallel
+ {
+ // CHECK-NEXT: OpenACCComputeConstruct {{.*}}parallel
+ // CHECK-NEXT: CompoundStmt
+#pragma acc parallel
+ {
+ // CHECK-NEXT: OpenACCComputeConstruct [[PAR_ADDR_UNINST:[0-9a-fx]+]] {{.*}}parallel
+ // CHECK-NEXT: CompoundStmt
+#pragma acc loop
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: [[PAR_ADDR_UNINST]]
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+ for(;;);
+ }
+ }
+
+ typename T::type array[5];
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl
+
+#pragma acc loop
+ for(auto x : array){}
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} <orphan>
+ // CHECK-NEXT: CXXForRangeStmt
+ // CHECK: CompoundStmt
+
+ // Instantiation:
+ // CHECK-NEXT: FunctionDecl{{.*}} TemplFunc 'void ()' implicit_instantiation
+ // CHECK-NEXT: TemplateArgument type 'S'
+ // CHECK-NEXT: RecordType{{.*}} 'S'
+ // CHECK-NEXT: CXXRecord{{.*}} 'S'
+ // CHECK-NEXT: CompoundStmt
+
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} <orphan>
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} used t 'typename S::type':'int'
+ // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 0
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: BinaryOperator{{.*}} 'bool' '<'
+ // CHECK-NEXT: ImplicitCastExpr{{.*}} 'typename S::type':'int' <LValueToRValue>
+ // CHECK-NEXT: DeclRefExpr {{.*}} 'typename S::type':'int' lvalue Var
+ // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 5
+ // CHECK-NEXT: UnaryOperator{{.*}} 'typename S::type':'int' lvalue prefix '++'
+ // CHECK-NEXT: DeclRefExpr {{.*}} 'typename S::type':'int' lvalue Var
+ // CHECK-NEXT: CompoundStmt
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} I 'typename S::type':'int'
+
+ // CHECK-NEXT: OpenACCComputeConstruct {{.*}}parallel
+ // CHECK-NEXT: CompoundStmt
+ //
+ // CHECK-NEXT: OpenACCComputeConstruct [[PAR_ADDR_INST:[0-9a-fx]+]] {{.*}}parallel
+ // CHECK-NEXT: CompoundStmt
+
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: [[PAR_ADDR_INST]]
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} <orphan>
+ // CHECK-NEXT: CXXForRangeStmt
+ // CHECK: CompoundStmt
+}
+
+struct S {
+ using type = int;
+};
+
+void use() {
+ TemplFunc<S>();
+}
+#endif
+
diff --git a/clang/test/SemaOpenACC/loop-loc-and-stmt.c b/clang/test/SemaOpenACC/loop-loc-and-stmt.c
new file mode 100644
index 0000000000000..36c6743f9843b
--- /dev/null
+++ b/clang/test/SemaOpenACC/loop-loc-and-stmt.c
@@ -0,0 +1,38 @@
+// RUN: %clang_cc1 %s -verify -fopenacc
+
+// expected-error at +1{{OpenACC construct 'loop' cannot be used here; it can only be used in a statement context}}
+#pragma acc loop
+
+// expected-error at +1{{OpenACC construct 'loop' cannot be used here; it can only be used in a statement context}}
+#pragma acc loop
+int foo;
+
+struct S {
+// expected-error at +1{{OpenACC construct 'loop' cannot be used here; it can only be used in a statement context}}
+#pragma acc loop
+ int i;
+};
+
+void func() {
+ // expected-error at +2{{expected expression}}
+#pragma acc loop
+ int foo;
+
+ // expected-error at +3{{OpenACC 'loop' construct can only be applied to a 'for' loop}}
+ // expected-note at +1{{'loop' construct is here}}
+#pragma acc loop
+ while(0);
+
+ // expected-error at +3{{OpenACC 'loop' construct can only be applied to a 'for' loop}}
+ // expected-note at +1{{'loop' construct is here}}
+#pragma acc loop
+ do{}while(0);
+
+ // expected-error at +3{{OpenACC 'loop' construct can only be applied to a 'for' loop}}
+ // expected-note at +1{{'loop' construct is here}}
+#pragma acc loop
+ {}
+
+#pragma acc loop
+ for(;;);
+}
diff --git a/clang/test/SemaOpenACC/loop-loc-and-stmt.cpp b/clang/test/SemaOpenACC/loop-loc-and-stmt.cpp
new file mode 100644
index 0000000000000..5d50145b7c882
--- /dev/null
+++ b/clang/test/SemaOpenACC/loop-loc-and-stmt.cpp
@@ -0,0 +1,80 @@
+// RUN: %clang_cc1 %s -verify -fopenacc
+//
+// expected-error at +1{{OpenACC construct 'loop' cannot be used here; it can only be used in a statement context}}
+#pragma acc loop
+
+// expected-error at +1{{OpenACC construct 'loop' cannot be used here; it can only be used in a statement context}}
+#pragma acc loop
+int foo;
+
+struct S {
+// expected-error at +1{{OpenACC construct 'loop' cannot be used here; it can only be used in a statement context}}
+#pragma acc loop
+ int i;
+
+ void mem_func() {
+ // expected-error at +3{{OpenACC 'loop' construct can only be applied to a 'for' loop}}
+ // expected-note at +1{{'loop' construct is here}}
+#pragma acc loop
+ int foo;
+
+ // expected-error at +3{{OpenACC 'loop' construct can only be applied to a 'for' loop}}
+ // expected-note at +1{{'loop' construct is here}}
+#pragma acc loop
+ while(0);
+
+ // expected-error at +3{{OpenACC 'loop' construct can only be applied to a 'for' loop}}
+ // expected-note at +1{{'loop' construct is here}}
+#pragma acc loop
+ do{}while(0);
+
+ // expected-error at +3{{OpenACC 'loop' construct can only be applied to a 'for' loop}}
+ // expected-note at +1{{'loop' construct is here}}
+#pragma acc loop
+ {}
+
+#pragma acc loop
+ for(;;);
+
+ int array[5];
+
+#pragma acc loop
+ for(auto X : array){}
+}
+};
+
+template<typename T>
+void templ_func() {
+ // expected-error at +3{{OpenACC 'loop' construct can only be applied to a 'for' loop}}
+ // expected-note at +1{{'loop' construct is here}}
+#pragma acc loop
+ int foo;
+
+ // expected-error at +3{{OpenACC 'loop' construct can only be applied to a 'for' loop}}
+ // expected-note at +1{{'loop' construct is here}}
+#pragma acc loop
+ while(T{});
+
+ // expected-error at +3{{OpenACC 'loop' construct can only be applied to a 'for' loop}}
+ // expected-note at +1{{'loop' construct is here}}
+#pragma acc loop
+ do{}while(0);
+
+ // expected-error at +3{{OpenACC 'loop' construct can only be applied to a 'for' loop}}
+ // expected-note at +1{{'loop' construct is here}}
+#pragma acc loop
+ {}
+
+#pragma acc loop
+ for(T i;;);
+
+ T array[5];
+
+#pragma acc loop
+ for(auto X : array){}
+}
+
+void use() {
+ templ_func<int>();
+}
+
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index 49ed60d990ca6..916e941cfbde1 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2170,6 +2170,7 @@ class EnqueueVisitor : public ConstStmtVisitor<EnqueueVisitor, void>,
void VisitRequiresExpr(const RequiresExpr *E);
void VisitCXXParenListInitExpr(const CXXParenListInitExpr *E);
void VisitOpenACCComputeConstruct(const OpenACCComputeConstruct *D);
+ void VisitOpenACCLoopConstruct(const OpenACCLoopConstruct *D);
void VisitOMPExecutableDirective(const OMPExecutableDirective *D);
void VisitOMPLoopBasedDirective(const OMPLoopBasedDirective *D);
void VisitOMPLoopDirective(const OMPLoopDirective *D);
@@ -3496,6 +3497,12 @@ void EnqueueVisitor::VisitOpenACCComputeConstruct(
EnqueueChildren(Clause);
}
+void EnqueueVisitor::VisitOpenACCLoopConstruct(const OpenACCLoopConstruct *C) {
+ EnqueueChildren(C);
+ for (auto *Clause : C->clauses())
+ EnqueueChildren(Clause);
+}
+
void EnqueueVisitor::VisitAnnotateAttr(const AnnotateAttr *A) {
EnqueueChildren(A);
}
@@ -6234,6 +6241,8 @@ CXString clang_getCursorKindSpelling(enum CXCursorKind Kind) {
return cxstring::createRef("ConceptDecl");
case CXCursor_OpenACCComputeConstruct:
return cxstring::createRef("OpenACCComputeConstruct");
+ case CXCursor_OpenACCLoopConstruct:
+ return cxstring::createRef("OpenACCLoopConstruct");
}
llvm_unreachable("Unhandled CXCursorKind");
diff --git a/clang/tools/libclang/CXCursor.cpp b/clang/tools/libclang/CXCursor.cpp
index 9325a16d2a848..38002052227cd 100644
--- a/clang/tools/libclang/CXCursor.cpp
+++ b/clang/tools/libclang/CXCursor.cpp
@@ -873,6 +873,9 @@ CXCursor cxcursor::MakeCXCursor(const Stmt *S, const Decl *Parent,
case Stmt::OpenACCComputeConstructClass:
K = CXCursor_OpenACCComputeConstruct;
break;
+ case Stmt::OpenACCLoopConstructClass:
+ K = CXCursor_OpenACCLoopConstruct;
+ break;
case Stmt::OMPTargetParallelGenericLoopDirectiveClass:
K = CXCursor_OMPTargetParallelGenericLoopDirective;
break;
>From 38de3cbcb63659c6a2d9caf05ce68b397b3a9250 Mon Sep 17 00:00:00 2001
From: erichkeane <ekeane at nvidia.com>
Date: Thu, 30 May 2024 06:34:59 -0700
Subject: [PATCH 2/3] Change visitor to use 'traverse' instead of 'visit'
What I really wanted was for the visitor searching for loops to stop
searching at the loop/compute construct, not to stop entirely. The
traverse functions allow us to do that, whereas the visit just 'visit'.
This patch switches to traverse so we stop correctly.
---
clang/lib/AST/StmtOpenACC.cpp | 9 +++++----
clang/test/SemaOpenACC/loop-ast.cpp | 18 ++++++++++++++++++
2 files changed, 23 insertions(+), 4 deletions(-)
diff --git a/clang/lib/AST/StmtOpenACC.cpp b/clang/lib/AST/StmtOpenACC.cpp
index e235449eb9959..61a6247aaf8b1 100644
--- a/clang/lib/AST/StmtOpenACC.cpp
+++ b/clang/lib/AST/StmtOpenACC.cpp
@@ -50,15 +50,16 @@ void OpenACCComputeConstruct::findAndSetChildLoops() {
LoopConstructFinder(OpenACCComputeConstruct *Construct)
: Construct(Construct) {}
- bool VisitOpenACCComputeConstruct(OpenACCComputeConstruct *C) {
+
+ bool TraverseOpenACCComputeConstruct(OpenACCComputeConstruct *C) {
// Stop searching if we find a compute construct.
- return false;
+ return true;
}
- bool VisitOpenACCLoopConstruct(OpenACCLoopConstruct *C) {
+ bool TraverseOpenACCLoopConstruct(OpenACCLoopConstruct *C) {
// Stop searching if we find a loop construct, after taking ownership of
// it.
C->setParentComputeConstruct(Construct);
- return false;
+ return true;
}
};
diff --git a/clang/test/SemaOpenACC/loop-ast.cpp b/clang/test/SemaOpenACC/loop-ast.cpp
index a55ca0b6a804c..292044f94267b 100644
--- a/clang/test/SemaOpenACC/loop-ast.cpp
+++ b/clang/test/SemaOpenACC/loop-ast.cpp
@@ -87,6 +87,16 @@ void TemplFunc() {
{
// CHECK-NEXT: OpenACCComputeConstruct [[PAR_ADDR_UNINST:[0-9a-fx]+]] {{.*}}parallel
// CHECK-NEXT: CompoundStmt
+#pragma acc loop
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: [[PAR_ADDR_UNINST]]
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+ for(;;);
+
#pragma acc loop
// CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: [[PAR_ADDR_UNINST]]
// CHECK-NEXT: ForStmt
@@ -146,6 +156,14 @@ void TemplFunc() {
// CHECK-NEXT: <<<NULL>>>
// CHECK-NEXT: NullStmt
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: [[PAR_ADDR_INST]]
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+
// CHECK-NEXT: DeclStmt
// CHECK-NEXT: VarDecl
// CHECK-NEXT: OpenACCLoopConstruct{{.*}} <orphan>
>From cc96a760b22bc0ef0760a89dc25c810835fb0fe3 Mon Sep 17 00:00:00 2001
From: erichkeane <ekeane at nvidia.com>
Date: Thu, 30 May 2024 06:39:36 -0700
Subject: [PATCH 3/3] Clang-format fix
---
clang/lib/AST/StmtOpenACC.cpp | 1 -
1 file changed, 1 deletion(-)
diff --git a/clang/lib/AST/StmtOpenACC.cpp b/clang/lib/AST/StmtOpenACC.cpp
index 61a6247aaf8b1..2d864a2885796 100644
--- a/clang/lib/AST/StmtOpenACC.cpp
+++ b/clang/lib/AST/StmtOpenACC.cpp
@@ -50,7 +50,6 @@ void OpenACCComputeConstruct::findAndSetChildLoops() {
LoopConstructFinder(OpenACCComputeConstruct *Construct)
: Construct(Construct) {}
-
bool TraverseOpenACCComputeConstruct(OpenACCComputeConstruct *C) {
// Stop searching if we find a compute construct.
return true;
More information about the cfe-commits
mailing list