[clang] [OpenACC] Add 'collapse' clause AST/basic Sema implementation (PR #109461)
Erich Keane via cfe-commits
cfe-commits at lists.llvm.org
Fri Sep 20 12:26:13 PDT 2024
https://github.com/erichkeane updated https://github.com/llvm/llvm-project/pull/109461
>From f25688a5f6242e4e16cad377fa281df665d0e38b Mon Sep 17 00:00:00 2001
From: erichkeane <ekeane at nvidia.com>
Date: Fri, 20 Sep 2024 09:38:32 -0700
Subject: [PATCH 1/2] [OpenACC] Add 'collapse' clause AST/basic Sema
implementation
The 'collapse' clause on a 'loop' construct is used to specify how many
nested loops are associated with the 'loop' construct. It takes an
optional 'force' tag, and an integer constant expression as arguments.
There are many other restrictions based on the contents of the loop/etc,
but those are implemented in followup patches, for now, this patch just
adds the AST node and does basic argument checking on the loop-count.
---
clang/include/clang/AST/OpenACCClause.h | 26 +++
.../clang/Basic/DiagnosticSemaKinds.td | 3 +
clang/include/clang/Basic/OpenACCClauses.def | 1 +
clang/include/clang/Sema/SemaOpenACC.h | 28 +++-
clang/lib/AST/OpenACCClause.cpp | 34 +++-
clang/lib/AST/StmtProfile.cpp | 6 +
clang/lib/AST/TextNodeDumper.cpp | 6 +
clang/lib/Parse/ParseOpenACC.cpp | 19 ++-
clang/lib/Sema/SemaOpenACC.cpp | 62 ++++++++
clang/lib/Sema/TreeTransform.h | 25 +++
clang/lib/Serialization/ASTReader.cpp | 8 +-
clang/lib/Serialization/ASTWriter.cpp | 8 +-
.../AST/ast-print-openacc-loop-construct.cpp | 25 +++
clang/test/ParserOpenACC/parse-clauses.c | 11 +-
clang/test/ParserOpenACC/parse-clauses.cpp | 18 ++-
.../compute-construct-device_type-clause.c | 3 +-
...p-construct-auto_seq_independent-clauses.c | 6 -
.../loop-construct-collapse-ast.cpp | 150 ++++++++++++++++++
.../loop-construct-collapse-clause.cpp | 117 ++++++++++++++
.../loop-construct-device_type-clause.c | 1 -
clang/tools/libclang/CIndex.cpp | 3 +
21 files changed, 532 insertions(+), 28 deletions(-)
create mode 100644 clang/test/SemaOpenACC/loop-construct-collapse-ast.cpp
create mode 100644 clang/test/SemaOpenACC/loop-construct-collapse-clause.cpp
diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index ea1ffbc7fd08b4..90f5b7fc9ab6f4 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -547,6 +547,32 @@ class OpenACCAsyncClause : public OpenACCClauseWithSingleIntExpr {
SourceLocation EndLoc);
};
+/// Represents a 'collapse' clause on a 'loop' construct. This clause takes an
+/// integer constant expression 'N' that represents how deep to collapse the
+/// construct. It also takes an optional 'force' tag that permits intervening
+/// code in the loops.
+class OpenACCCollapseClause : public OpenACCClauseWithSingleIntExpr {
+ bool HasForce = false;
+
+ OpenACCCollapseClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
+ bool HasForce, Expr *LoopCount, SourceLocation EndLoc);
+
+public:
+ const Expr *getLoopCount() const { return getIntExpr(); }
+ Expr *getLoopCount() { return getIntExpr(); }
+
+ bool hasForce() const { return HasForce; }
+
+ static bool classof(const OpenACCClause *C) {
+ return C->getClauseKind() == OpenACCClauseKind::Collapse;
+ }
+
+ static OpenACCCollapseClause *Create(const ASTContext &C,
+ SourceLocation BeginLoc,
+ SourceLocation LParenLoc, bool HasForce,
+ Expr *LoopCount, SourceLocation EndLoc);
+};
+
/// Represents a clause with one or more 'var' objects, represented as an expr,
/// as its arguments. Var-list is expected to be stored in trailing storage.
/// For now, we're just storing the original expression in its entirety, unlike
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index ba813af960af6f..fb60477bc0c2e6 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12606,6 +12606,9 @@ def note_acc_construct_here : Note<"'%0' construct is here">;
def err_acc_loop_spec_conflict
: Error<"OpenACC clause '%0' on '%1' construct conflicts with previous "
"data dependence clause">;
+def err_acc_collapse_loop_count
+ : Error<"OpenACC 'collapse' clause loop count must be a %select{constant "
+ "expression|positive integer value, evaluated to %1}0">;
// AMDGCN builtins diagnostics
def err_amdgcn_global_load_lds_size_invalid_value : Error<"invalid size value">;
diff --git a/clang/include/clang/Basic/OpenACCClauses.def b/clang/include/clang/Basic/OpenACCClauses.def
index 85f4859925f0bc..19cdfe7672133b 100644
--- a/clang/include/clang/Basic/OpenACCClauses.def
+++ b/clang/include/clang/Basic/OpenACCClauses.def
@@ -24,6 +24,7 @@
VISIT_CLAUSE(Auto)
VISIT_CLAUSE(Async)
VISIT_CLAUSE(Attach)
+VISIT_CLAUSE(Collapse)
VISIT_CLAUSE(Copy)
CLAUSE_ALIAS(PCopy, Copy, true)
CLAUSE_ALIAS(PresentOrCopy, Copy, true)
diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index 0ca76842e5f902..b4d54d95953046 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -87,9 +87,14 @@ class SemaOpenACC : public SemaBase {
SmallVector<Expr *> VarList;
};
+ struct CollapseDetails {
+ bool IsForce;
+ Expr *LoopCount;
+ };
+
std::variant<std::monostate, DefaultDetails, ConditionDetails,
IntExprDetails, VarListDetails, WaitDetails, DeviceTypeDetails,
- ReductionDetails>
+ ReductionDetails, CollapseDetails>
Details = std::monostate{};
public:
@@ -246,6 +251,19 @@ class SemaOpenACC : public SemaBase {
return std::get<VarListDetails>(Details).IsZero;
}
+ bool isForce() const {
+ assert(ClauseKind == OpenACCClauseKind::Collapse &&
+ "Only 'collapse' has a force tag");
+ return std::get<CollapseDetails>(Details).IsForce;
+ }
+
+ Expr *getLoopCount() const {
+ assert(ClauseKind == OpenACCClauseKind::Collapse &&
+ "Only 'collapse' has a loop count");
+ return std::get<CollapseDetails>(Details).LoopCount;
+ }
+
+
ArrayRef<DeviceTypeArgument> getDeviceTypeArchitectures() const {
assert((ClauseKind == OpenACCClauseKind::DeviceType ||
ClauseKind == OpenACCClauseKind::DType) &&
@@ -384,6 +402,12 @@ class SemaOpenACC : public SemaBase {
"Only 'device_type'/'dtype' has a device-type-arg list");
Details = DeviceTypeDetails{std::move(Archs)};
}
+
+ void setCollapseDetails(bool IsForce, Expr *LoopCount) {
+ assert(ClauseKind == OpenACCClauseKind::Collapse &&
+ "Only 'collapse' has collapse details");
+ Details = CollapseDetails{IsForce, LoopCount};
+ }
};
SemaOpenACC(Sema &S);
@@ -448,6 +472,8 @@ class SemaOpenACC : public SemaBase {
Expr *LowerBound,
SourceLocation ColonLocFirst, Expr *Length,
SourceLocation RBLoc);
+ /// Checks the loop depth value for a collapse clause.
+ ExprResult CheckCollapseLoopCount(Expr *LoopCount);
/// Helper type for the registration/assignment of constructs that need to
/// 'know' about their parent constructs and hold a reference to them, such as
diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp
index 95089a9b79e267..fc8a4c28dc3fe4 100644
--- a/clang/lib/AST/OpenACCClause.cpp
+++ b/clang/lib/AST/OpenACCClause.cpp
@@ -43,7 +43,7 @@ bool OpenACCClauseWithCondition::classof(const OpenACCClause *C) {
bool OpenACCClauseWithSingleIntExpr::classof(const OpenACCClause *C) {
return OpenACCNumWorkersClause::classof(C) ||
OpenACCVectorLengthClause::classof(C) ||
- OpenACCAsyncClause::classof(C);
+ OpenACCCollapseClause::classof(C) || OpenACCAsyncClause::classof(C);
}
OpenACCDefaultClause *OpenACCDefaultClause::Create(const ASTContext &C,
OpenACCDefaultClauseKind K,
@@ -134,6 +134,30 @@ OpenACCNumWorkersClause::Create(const ASTContext &C, SourceLocation BeginLoc,
OpenACCNumWorkersClause(BeginLoc, LParenLoc, IntExpr, EndLoc);
}
+OpenACCCollapseClause::OpenACCCollapseClause(SourceLocation BeginLoc,
+ SourceLocation LParenLoc,
+ bool HasForce, Expr *LoopCount,
+ SourceLocation EndLoc)
+ : OpenACCClauseWithSingleIntExpr(OpenACCClauseKind::Collapse, BeginLoc,
+ LParenLoc, LoopCount, EndLoc),
+ HasForce(HasForce) {
+ assert(LoopCount && "LoopCount required");
+}
+
+OpenACCCollapseClause *
+OpenACCCollapseClause::Create(const ASTContext &C, SourceLocation BeginLoc,
+ SourceLocation LParenLoc, bool HasForce,
+ Expr *LoopCount, SourceLocation EndLoc) {
+ assert(LoopCount &&
+ (LoopCount->isInstantiationDependent() ||
+ (LoopCount->isIntegerConstantExpr(C))) &&
+ "Loop count not constant expression");
+ void *Mem =
+ C.Allocate(sizeof(OpenACCCollapseClause), alignof(OpenACCCollapseClause));
+ return new (Mem)
+ OpenACCCollapseClause(BeginLoc, LParenLoc, HasForce, LoopCount, EndLoc);
+}
+
OpenACCVectorLengthClause::OpenACCVectorLengthClause(SourceLocation BeginLoc,
SourceLocation LParenLoc,
Expr *IntExpr,
@@ -550,3 +574,11 @@ void OpenACCClausePrinter::VisitIndependentClause(
void OpenACCClausePrinter::VisitSeqClause(const OpenACCSeqClause &C) {
OS << "seq";
}
+
+void OpenACCClausePrinter::VisitCollapseClause(const OpenACCCollapseClause &C) {
+ OS << "collapse(";
+ if (C.hasForce())
+ OS << "force:";
+ printExpr(C.getLoopCount());
+ OS << ")";
+}
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index ad4281986f668e..c3812844ab8a31 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -2558,6 +2558,12 @@ void OpenACCClauseProfiler::VisitNumWorkersClause(
Profiler.VisitStmt(Clause.getIntExpr());
}
+void OpenACCClauseProfiler::VisitCollapseClause(
+ const OpenACCCollapseClause &Clause) {
+ assert(Clause.getLoopCount() && "collapse clause requires a valid int expr");
+ Profiler.VisitStmt(Clause.getLoopCount());
+}
+
void OpenACCClauseProfiler::VisitPrivateClause(
const OpenACCPrivateClause &Clause) {
for (auto *E : Clause.getVarList())
diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp
index 3c51c746471829..8a74159c7c93e5 100644
--- a/clang/lib/AST/TextNodeDumper.cpp
+++ b/clang/lib/AST/TextNodeDumper.cpp
@@ -419,6 +419,12 @@ void TextNodeDumper::Visit(const OpenACCClause *C) {
// but print 'clause' here so it is clear what is happening from the dump.
OS << " clause";
break;
+ case OpenACCClauseKind::Collapse:
+ OS << " clause";
+ if (cast<OpenACCCollapseClause>(C)->hasForce())
+ OS << ": force";
+ break;
+
case OpenACCClauseKind::CopyIn:
case OpenACCClauseKind::PCopyIn:
case OpenACCClauseKind::PresentOrCopyIn:
diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
index 0261e8ea3c9b76..e66abd6873794e 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -976,14 +976,25 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
/*IsReadOnly=*/false, /*IsZero=*/false);
break;
case OpenACCClauseKind::Collapse: {
- tryParseAndConsumeSpecialTokenKind(*this, OpenACCSpecialTokenKind::Force,
- ClauseKind);
- ExprResult NumLoops =
+ bool HasForce = tryParseAndConsumeSpecialTokenKind(
+ *this, OpenACCSpecialTokenKind::Force, ClauseKind);
+ ExprResult LoopCount =
getActions().CorrectDelayedTyposInExpr(ParseConstantExpression());
- if (NumLoops.isInvalid()) {
+ if (LoopCount.isInvalid()) {
Parens.skipToEnd();
return OpenACCCanContinue();
}
+
+ LoopCount = getActions().OpenACC().ActOnIntExpr(
+ OpenACCDirectiveKind::Invalid, ClauseKind,
+ LoopCount.get()->getBeginLoc(), LoopCount.get());
+
+ if (LoopCount.isInvalid()) {
+ Parens.skipToEnd();
+ return OpenACCCanContinue();
+ }
+
+ ParsedClause.setCollapseDetails(HasForce, LoopCount.get());
break;
}
case OpenACCClauseKind::Bind: {
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index ecbcc19413dc61..6a4054b53fcb43 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -343,6 +343,18 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
return false;
}
+ case OpenACCClauseKind::Collapse: {
+ switch (DirectiveKind) {
+ case OpenACCDirectiveKind::Loop:
+ case OpenACCDirectiveKind::ParallelLoop:
+ case OpenACCDirectiveKind::SerialLoop:
+ case OpenACCDirectiveKind::KernelsLoop:
+ return true;
+ default:
+ return false;
+ }
+ }
+
default:
// Do nothing so we can go to the 'unimplemented' diagnostic instead.
return true;
@@ -1037,6 +1049,26 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitReductionClause(
ValidVars, Clause.getEndLoc());
}
+OpenACCClause *SemaOpenACCClauseVisitor::VisitCollapseClause(
+ SemaOpenACC::OpenACCParsedClause &Clause) {
+ // Duplicates here are not really sensible. We could possible permit
+ // multiples if they all had the same value, but there isn't really a good
+ // reason to do so. Also, this simplifies the suppression of duplicates, in
+ // that we know if we 'find' one after instantiation, that it is the same
+ // clause, which simplifies instantiation/checking/etc.
+ if (checkAlreadyHasClauseOfKind(SemaRef, ExistingClauses, Clause))
+ return nullptr;
+
+ ExprResult LoopCount = SemaRef.CheckCollapseLoopCount(Clause.getLoopCount());
+
+ if (!LoopCount.isUsable())
+ return nullptr;
+
+ return OpenACCCollapseClause::Create(
+ Ctx, Clause.getBeginLoc(), Clause.getLParenLoc(),
+ Clause.isForce(), LoopCount.get(), Clause.getEndLoc());
+}
+
} // namespace
SemaOpenACC::SemaOpenACC(Sema &S) : SemaBase(S) {}
@@ -1273,6 +1305,9 @@ ExprResult SemaOpenACC::ActOnIntExpr(OpenACCDirectiveKind DK,
}
} IntExprDiagnoser(DK, CK, IntExpr);
+ if (!IntExpr)
+ return ExprError();
+
ExprResult IntExprResult = SemaRef.PerformContextualImplicitConversion(
Loc, IntExpr, IntExprDiagnoser);
if (IntExprResult.isInvalid())
@@ -1583,6 +1618,33 @@ ExprResult SemaOpenACC::ActOnArraySectionExpr(Expr *Base, SourceLocation LBLoc,
OK_Ordinary, ColonLoc, RBLoc);
}
+ExprResult SemaOpenACC::CheckCollapseLoopCount(Expr *LoopCount) {
+ if (!LoopCount)
+ return ExprError();
+
+ assert((LoopCount->isInstantiationDependent() ||
+ LoopCount->getType()->isIntegerType()) &&
+ "Loop argument non integer?");
+
+ // If this is dependent, there really isn't anything we can check.
+ if (LoopCount->isInstantiationDependent())
+ return ExprResult{LoopCount};
+
+ std::optional<llvm::APSInt> ICE =
+ LoopCount->getIntegerConstantExpr(getASTContext());
+
+ // OpenACC 3.3: 2.9.1
+ // The argument to the collapse clause must be a constant positive integer
+ // expression.
+ if (!ICE || *ICE <= 0) {
+ Diag(LoopCount->getBeginLoc(), diag::err_acc_collapse_loop_count)
+ << ICE.has_value() << ICE.value_or(llvm::APSInt{}).getExtValue();
+ return ExprError();
+ }
+
+ return ExprResult{LoopCount};
+}
+
bool SemaOpenACC::ActOnStartStmtDirective(OpenACCDirectiveKind K,
SourceLocation StartLoc) {
SemaRef.DiscardCleanupsInEvaluationContext();
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index 95ded5e59a9fa7..3dba8e53669e8d 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -11821,6 +11821,31 @@ void OpenACCClauseTransform<Derived>::VisitReductionClause(
ParsedClause.getLParenLoc(), C.getReductionOp(), ValidVars,
ParsedClause.getEndLoc());
}
+
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitCollapseClause(
+ const OpenACCCollapseClause &C) {
+ Expr *LoopCount = const_cast<Expr *>(C.getLoopCount());
+ assert(LoopCount && "collapse clause constructed with invalid loop count");
+
+ ExprResult NewLoopCount = Self.TransformExpr(LoopCount);
+
+ NewLoopCount = Self.getSema().OpenACC().ActOnIntExpr(
+ OpenACCDirectiveKind::Invalid, ParsedClause.getClauseKind(),
+ NewLoopCount.get()->getBeginLoc(), NewLoopCount.get());
+
+ NewLoopCount =
+ Self.getSema().OpenACC().CheckCollapseLoopCount(NewLoopCount.get());
+
+ if (!NewLoopCount.isUsable())
+ return;
+
+ ParsedClause.setCollapseDetails(C.hasForce(), NewLoopCount.get());
+ NewClause = OpenACCCollapseClause::Create(
+ Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
+ ParsedClause.getLParenLoc(), ParsedClause.isForce(),
+ ParsedClause.getLoopCount(), ParsedClause.getEndLoc());
+}
} // namespace
template <typename Derived>
OpenACCClause *TreeTransform<Derived>::TransformOpenACCClause(
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index 7efcc81e194d95..7a0a6d81b21a58 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -12269,6 +12269,13 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
return OpenACCIndependentClause::Create(getContext(), BeginLoc, EndLoc);
case OpenACCClauseKind::Auto:
return OpenACCAutoClause::Create(getContext(), BeginLoc, EndLoc);
+ case OpenACCClauseKind::Collapse: {
+ SourceLocation LParenLoc = readSourceLocation();
+ bool HasForce = readBool();
+ Expr *LoopCount = readSubExpr();
+ return OpenACCCollapseClause::Create(getContext(), BeginLoc, LParenLoc,
+ HasForce, LoopCount, EndLoc);
+ }
case OpenACCClauseKind::Finalize:
case OpenACCClauseKind::IfPresent:
@@ -12282,7 +12289,6 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
case OpenACCClauseKind::DeviceResident:
case OpenACCClauseKind::Host:
case OpenACCClauseKind::Link:
- case OpenACCClauseKind::Collapse:
case OpenACCClauseKind::Bind:
case OpenACCClauseKind::DeviceNum:
case OpenACCClauseKind::DefaultAsync:
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index 008bf571f847dc..2da0c94c9e4cbc 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -8126,6 +8126,13 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
// Nothing to do here, there is no additional information beyond the
// begin/end loc and clause kind.
return;
+ case OpenACCClauseKind::Collapse: {
+ const auto *CC = cast<OpenACCCollapseClause>(C);
+ writeSourceLocation(CC->getLParenLoc());
+ writeBool(CC->hasForce());
+ AddStmt(const_cast<Expr *>(CC->getLoopCount()));
+ return;
+ }
case OpenACCClauseKind::Finalize:
case OpenACCClauseKind::IfPresent:
@@ -8139,7 +8146,6 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
case OpenACCClauseKind::DeviceResident:
case OpenACCClauseKind::Host:
case OpenACCClauseKind::Link:
- case OpenACCClauseKind::Collapse:
case OpenACCClauseKind::Bind:
case OpenACCClauseKind::DeviceNum:
case OpenACCClauseKind::DefaultAsync:
diff --git a/clang/test/AST/ast-print-openacc-loop-construct.cpp b/clang/test/AST/ast-print-openacc-loop-construct.cpp
index cde302a66f3af7..ae1f7964f019eb 100644
--- a/clang/test/AST/ast-print-openacc-loop-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-loop-construct.cpp
@@ -57,4 +57,29 @@ void foo() {
// CHECK-NEXT: ;
#pragma acc loop private(i, array[1], array, array[1:2])
for(;;);
+
+// CHECK: #pragma acc loop collapse(1)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc loop collapse(1)
+ for(;;);
+// CHECK: #pragma acc loop collapse(force:1)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc loop collapse(force:1)
+ for(;;);
+// CHECK: #pragma acc loop collapse(2)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc loop collapse(2)
+ for(;;)
+ for(;;);
+// CHECK: #pragma acc loop collapse(force:2)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc loop collapse(force:2)
+ for(;;)
+ for(;;);
}
diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c
index 15c4554a31922a..6c9ce4ad5e1969 100644
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@@ -105,17 +105,14 @@ void func() {
#pragma acc loop collapse(force:)
for(;;){}
- // 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)
+ // expected-error at +1{{invalid tag 'unknown' on 'collapse' clause}}
+#pragma acc loop collapse(unknown:1)
for(;;){}
- // expected-warning at +1{{OpenACC clause 'collapse' not yet implemented, clause ignored}}
-#pragma acc loop collapse(force:5)
+#pragma acc loop collapse(force:1)
for(;;){}
- // expected-warning at +1{{OpenACC clause 'collapse' not yet implemented, clause ignored}}
-#pragma acc loop collapse(5)
+#pragma acc loop collapse(1)
for(;;){}
// expected-error at +2{{expected ')'}}
diff --git a/clang/test/ParserOpenACC/parse-clauses.cpp b/clang/test/ParserOpenACC/parse-clauses.cpp
index b7e252e892beab..9613530db77ddc 100644
--- a/clang/test/ParserOpenACC/parse-clauses.cpp
+++ b/clang/test/ParserOpenACC/parse-clauses.cpp
@@ -2,13 +2,23 @@
template<unsigned I, typename T>
void templ() {
- // expected-warning at +1{{OpenACC clause 'collapse' not yet implemented, clause ignored}}
#pragma acc loop collapse(I)
- for(;;){}
+ for(;;)
+ for(;;)
+ for(;;)
+ for(;;)
+ for(;;)
+ for(;;)
+ for(;;);
- // expected-warning at +1{{OpenACC clause 'collapse' not yet implemented, clause ignored}}
#pragma acc loop collapse(T::value)
- for(;;){}
+ for(;;)
+ for(;;)
+ for(;;)
+ for(;;)
+ for(;;)
+ for(;;)
+ for(;;);
#pragma acc parallel vector_length(T::value)
for(;;){}
diff --git a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
index b300abe577801c..26f0315fb86f1a 100644
--- a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
@@ -188,8 +188,7 @@ void uses() {
// expected-note at +1{{previous clause is here}}
#pragma acc serial device_type(*) reduction(+:Var)
while(1);
- // expected-error at +2{{OpenACC clause 'collapse' may not follow a 'device_type' clause in a compute construct}}
- // expected-note at +1{{previous clause is here}}
+ // expected-error at +1{{OpenACC 'collapse' clause is not valid on 'kernels' directive}}
#pragma acc kernels device_type(*) collapse(1)
while(1);
// expected-error at +2{{OpenACC clause 'bind' may not follow a 'device_type' clause in a compute construct}}
diff --git a/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c b/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
index ac61976ff620d8..3212c19d089fc9 100644
--- a/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
+++ b/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
@@ -138,7 +138,6 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'reduction' not yet implemented}}
#pragma acc loop auto reduction(+:Var)
for(;;);
- // expected-warning at +1{{OpenACC clause 'collapse' not yet implemented}}
#pragma acc loop auto collapse(1)
for(;;);
// expected-warning at +1{{OpenACC clause 'bind' not yet implemented}}
@@ -277,7 +276,6 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'reduction' not yet implemented}}
#pragma acc loop reduction(+:Var) auto
for(;;);
- // expected-warning at +1{{OpenACC clause 'collapse' not yet implemented}}
#pragma acc loop collapse(1) auto
for(;;);
// expected-warning at +1{{OpenACC clause 'bind' not yet implemented}}
@@ -417,7 +415,6 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'reduction' not yet implemented}}
#pragma acc loop independent reduction(+:Var)
for(;;);
- // expected-warning at +1{{OpenACC clause 'collapse' not yet implemented}}
#pragma acc loop independent collapse(1)
for(;;);
// expected-warning at +1{{OpenACC clause 'bind' not yet implemented}}
@@ -556,7 +553,6 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'reduction' not yet implemented}}
#pragma acc loop reduction(+:Var) independent
for(;;);
- // expected-warning at +1{{OpenACC clause 'collapse' not yet implemented}}
#pragma acc loop collapse(1) independent
for(;;);
// expected-warning at +1{{OpenACC clause 'bind' not yet implemented}}
@@ -705,7 +701,6 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'reduction' not yet implemented}}
#pragma acc loop seq reduction(+:Var)
for(;;);
- // expected-warning at +1{{OpenACC clause 'collapse' not yet implemented}}
#pragma acc loop seq collapse(1)
for(;;);
// expected-warning at +1{{OpenACC clause 'bind' not yet implemented}}
@@ -853,7 +848,6 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'reduction' not yet implemented}}
#pragma acc loop reduction(+:Var) seq
for(;;);
- // expected-warning at +1{{OpenACC clause 'collapse' not yet implemented}}
#pragma acc loop collapse(1) seq
for(;;);
// expected-warning at +1{{OpenACC clause 'bind' not yet implemented}}
diff --git a/clang/test/SemaOpenACC/loop-construct-collapse-ast.cpp b/clang/test/SemaOpenACC/loop-construct-collapse-ast.cpp
new file mode 100644
index 00000000000000..4048c450f8fa1d
--- /dev/null
+++ b/clang/test/SemaOpenACC/loop-construct-collapse-ast.cpp
@@ -0,0 +1,150 @@
+// 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
+
+struct S {
+ constexpr S(){};
+ constexpr operator auto() {return 1;}
+};
+
+void NormalUses() {
+ // CHECK: FunctionDecl{{.*}}NormalUses
+ // CHECK-NEXT: CompoundStmt
+
+#pragma acc loop collapse(1)
+ for(;;);
+ // CHECK-NEXT: OpenACCLoopConstruct
+ // CHECK-NEXT: collapse clause
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+
+#pragma acc loop collapse(force:S{})
+ for(;;);
+ // CHECK-NEXT: OpenACCLoopConstruct
+ // CHECK-NEXT: collapse clause
+ // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <UserDefinedConversion>
+ // CHECK-NEXT: CXXMemberCallExpr{{.*}} 'int'
+ // CHECK-NEXT: MemberExpr{{.*}} .operator auto
+ // CHECK-NEXT: MaterializeTemporaryExpr{{.*}}'S' lvalue
+ // CHECK-NEXT: CXXTemporaryObjectExpr{{.*}}'S' 'void ()' list
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+}
+
+template<typename T, unsigned Value>
+void TemplUses() {
+ // CHECK: FunctionTemplateDecl{{.*}}TemplUses
+ // CHECK-NEXT: TemplateTypeParmDecl {{.*}} referenced typename depth 0 index 0 T
+ // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} referenced 'unsigned int' depth 0 index 1 Value
+ // CHECK-NEXT: FunctionDecl{{.*}} TemplUses 'void ()'
+ // CHECK-NEXT: CompoundStmt
+
+#pragma acc loop collapse(Value)
+ for(;;)
+ for(;;);
+ // CHECK-NEXT: OpenACCLoopConstruct
+ // CHECK-NEXT: collapse clause
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'unsigned int' NonTypeTemplateParm{{.*}} 'Value'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+
+#pragma acc loop collapse(force:T{} + S{})
+ for(;;)
+ for(;;);
+ // CHECK-NEXT: OpenACCLoopConstruct
+ // CHECK-NEXT: collapse clause
+ // CHECK-NEXT: BinaryOperator {{.*}}'+'
+ // CHECK-NEXT: CXXUnresolvedConstructExpr{{.*}}'T' 'T' list
+ // CHECK-NEXT: InitListExpr
+ // CHECK-NEXT: CXXTemporaryObjectExpr{{.*}}'S' 'void ()' list
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+
+ // Instantiation:
+ // CHECK-NEXT: FunctionDecl{{.*}} used TemplUses 'void ()' implicit_instantiation
+ // CHECK-NEXT: TemplateArgument type 'S'
+ // CHECK-NEXT: RecordType{{.*}} 'S'
+ // CHECK-NEXT: CXXRecord{{.*}} 'S'
+ // CHECK-NEXT: TemplateArgument integral '2U'
+ // CHECK-NEXT: CompoundStmt
+
+ // CHECK-NEXT: OpenACCLoopConstruct
+ // CHECK-NEXT: collapse clause
+ // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}} 'unsigned int'
+ // CHECK-NEXT: NonTypeTemplateParmDecl
+ // CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned int' 2
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+
+ // CHECK-NEXT: OpenACCLoopConstruct
+ // CHECK-NEXT: collapse clause
+ // CHECK-NEXT: BinaryOperator {{.*}}'+'
+ // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <UserDefinedConversion>
+ // CHECK-NEXT: CXXMemberCallExpr{{.*}} 'int'
+ // CHECK-NEXT: MemberExpr{{.*}} .operator auto
+ // CHECK-NEXT: MaterializeTemporaryExpr{{.*}}'S' lvalue
+ // CHECK-NEXT: CXXTemporaryObjectExpr{{.*}}'S' 'void ()' list
+ // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <UserDefinedConversion>
+ // CHECK-NEXT: CXXMemberCallExpr{{.*}} 'int'
+ // CHECK-NEXT: MemberExpr{{.*}} .operator auto
+ // CHECK-NEXT: MaterializeTemporaryExpr{{.*}}'S' lvalue
+ // CHECK-NEXT: CXXTemporaryObjectExpr{{.*}}'S' 'void ()' list
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+
+}
+
+void Inst() {
+ TemplUses<S, 2>();
+}
+
+#endif // PCH_HELPER
diff --git a/clang/test/SemaOpenACC/loop-construct-collapse-clause.cpp b/clang/test/SemaOpenACC/loop-construct-collapse-clause.cpp
new file mode 100644
index 00000000000000..9c1e577773e8f8
--- /dev/null
+++ b/clang/test/SemaOpenACC/loop-construct-collapse-clause.cpp
@@ -0,0 +1,117 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+
+void only_for_loops() {
+ // 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 collapse(1)
+ while(true);
+
+ // 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 collapse(1)
+ do{}while(true);
+
+}
+
+void only_one_on_loop() {
+ // expected-error at +2{{OpenACC 'collapse' clause cannot appear more than once on a 'loop' directive}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc loop collapse(1) collapse(1)
+ for(;;);
+}
+
+constexpr int three() { return 3; }
+constexpr int one() { return 1; }
+constexpr int neg() { return -1; }
+constexpr int zero() { return 0; }
+
+struct NotConstexpr {
+ constexpr NotConstexpr(){};
+
+ operator int(){ return 1; }
+};
+struct ConvertsNegative {
+ constexpr ConvertsNegative(){};
+
+ constexpr operator int(){ return -1; }
+};
+struct ConvertsOne{
+ constexpr ConvertsOne(){};
+
+ constexpr operator int(){ return 1; }
+};
+
+struct ConvertsThree{
+ constexpr ConvertsThree(){};
+
+ constexpr operator int(){ return 3; }
+};
+
+template <typename T, int Val>
+void negative_constexpr_templ() {
+ // expected-error at +3 2{{OpenACC 'collapse' clause loop count must be a positive integer value, evaluated to 0}}
+ // expected-note@#NCETN1{{in instantiation of function template specialization 'negative_constexpr_templ<int, -1>'}}
+ // expected-note@#NCET1{{in instantiation of function template specialization 'negative_constexpr_templ<int, 1>'}}
+#pragma acc loop collapse(T{})
+ for(;;)
+ for(;;);
+
+ // expected-error at +1{{OpenACC 'collapse' clause loop count must be a positive integer value, evaluated to -1}}
+#pragma acc loop collapse(Val)
+ for(;;)
+ for(;;);
+}
+
+void negative_constexpr(int i) {
+#pragma acc loop collapse(2)
+ for(;;)
+ for(;;);
+
+#pragma acc loop collapse(1)
+ for(;;)
+ for(;;);
+
+ // expected-error at +1{{OpenACC 'collapse' clause loop count must be a positive integer value, evaluated to 0}}
+#pragma acc loop collapse(0)
+ for(;;)
+ for(;;);
+
+ // expected-error at +1{{OpenACC 'collapse' clause loop count must be a positive integer value, evaluated to -1}}
+#pragma acc loop collapse(-1)
+ for(;;)
+ for(;;);
+
+#pragma acc loop collapse(one())
+ for(;;)
+ for(;;);
+
+ // expected-error at +1{{OpenACC 'collapse' clause loop count must be a positive integer value, evaluated to 0}}
+#pragma acc loop collapse(zero())
+ for(;;)
+ for(;;);
+
+ // expected-error at +1{{OpenACC 'collapse' clause loop count must be a positive integer value, evaluated to -1}}
+#pragma acc loop collapse(neg())
+ for(;;)
+ for(;;);
+
+ // expected-error at +1{{OpenACC 'collapse' clause loop count must be a constant expression}}
+#pragma acc loop collapse(NotConstexpr{})
+ for(;;)
+ for(;;);
+
+ // expected-error at +1{{OpenACC 'collapse' clause loop count must be a positive integer value, evaluated to -1}}
+#pragma acc loop collapse(ConvertsNegative{})
+ for(;;)
+ for(;;);
+
+#pragma acc loop collapse(ConvertsOne{})
+ for(;;)
+ for(;;);
+
+ negative_constexpr_templ<int, -1>(); // #NCETN1
+
+ negative_constexpr_templ<int, 1>(); // #NCET1
+}
+
diff --git a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
index 520ba45aaebf4f..47c9239f4f0e96 100644
--- a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
@@ -162,7 +162,6 @@ void uses() {
// expected-note at +1{{previous clause is here}}
#pragma acc loop device_type(*) reduction(+:Var)
for(;;);
- // expected-warning at +1{{OpenACC clause 'collapse' not yet implemented, clause ignored}}
#pragma acc loop device_type(*) collapse(1)
for(;;);
// expected-error at +2{{OpenACC clause 'bind' may not follow a 'device_type' clause in a 'loop' construct}}
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index e821c5e4c588b6..83a015b8f3ca71 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2893,6 +2893,9 @@ void OpenACCClauseEnqueue::VisitAutoClause(const OpenACCAutoClause &C) {}
void OpenACCClauseEnqueue::VisitIndependentClause(
const OpenACCIndependentClause &C) {}
void OpenACCClauseEnqueue::VisitSeqClause(const OpenACCSeqClause &C) {}
+void OpenACCClauseEnqueue::VisitCollapseClause(const OpenACCCollapseClause &C) {
+ Visitor.AddStmt(C.getLoopCount());
+}
} // namespace
void EnqueueVisitor::EnqueueChildren(const OpenACCClause *C) {
>From de75b8459a6340ab9bc4f2910f2d5204a9cea89a Mon Sep 17 00:00:00 2001
From: erichkeane <ekeane at nvidia.com>
Date: Fri, 20 Sep 2024 12:26:02 -0700
Subject: [PATCH 2/2] Clang-format
---
clang/include/clang/Sema/SemaOpenACC.h | 1 -
clang/lib/Sema/SemaOpenACC.cpp | 6 +++---
2 files changed, 3 insertions(+), 4 deletions(-)
diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index b4d54d95953046..839fdb79cd0ac8 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -263,7 +263,6 @@ class SemaOpenACC : public SemaBase {
return std::get<CollapseDetails>(Details).LoopCount;
}
-
ArrayRef<DeviceTypeArgument> getDeviceTypeArchitectures() const {
assert((ClauseKind == OpenACCClauseKind::DeviceType ||
ClauseKind == OpenACCClauseKind::DType) &&
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 6a4054b53fcb43..18d02402efe04f 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -1064,9 +1064,9 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitCollapseClause(
if (!LoopCount.isUsable())
return nullptr;
- return OpenACCCollapseClause::Create(
- Ctx, Clause.getBeginLoc(), Clause.getLParenLoc(),
- Clause.isForce(), LoopCount.get(), Clause.getEndLoc());
+ return OpenACCCollapseClause::Create(Ctx, Clause.getBeginLoc(),
+ Clause.getLParenLoc(), Clause.isForce(),
+ LoopCount.get(), Clause.getEndLoc());
}
} // namespace
More information about the cfe-commits
mailing list