[clang] d412cea - [OpenACC] Implement 'tile' attribute AST (#110999)
via cfe-commits
cfe-commits at lists.llvm.org
Thu Oct 3 08:34:47 PDT 2024
Author: Erich Keane
Date: 2024-10-03T08:34:43-07:00
New Revision: d412cea8c4f26f451aee46641e384e8df62a5904
URL: https://github.com/llvm/llvm-project/commit/d412cea8c4f26f451aee46641e384e8df62a5904
DIFF: https://github.com/llvm/llvm-project/commit/d412cea8c4f26f451aee46641e384e8df62a5904.diff
LOG: [OpenACC] Implement 'tile' attribute AST (#110999)
The 'tile' clause shares quite a bit of the rules with 'collapse', so a
followup patch will add those tests/behaviors. This patch deals with
adding the AST node.
The 'tile' clause takes a series of integer constant expressions, or *.
The asterisk is now represented by a new OpenACCAsteriskSizeExpr node,
else this clause is very similar to others.
Added:
clang/test/SemaOpenACC/loop-construct-tile-ast.cpp
clang/test/SemaOpenACC/loop-construct-tile-clause.cpp
Modified:
clang/include/clang/AST/ComputeDependence.h
clang/include/clang/AST/Expr.h
clang/include/clang/AST/JSONNodeDumper.h
clang/include/clang/AST/OpenACCClause.h
clang/include/clang/AST/RecursiveASTVisitor.h
clang/include/clang/AST/TextNodeDumper.h
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/include/clang/Basic/OpenACCClauses.def
clang/include/clang/Basic/StmtNodes.td
clang/include/clang/Parse/Parser.h
clang/include/clang/Sema/SemaOpenACC.h
clang/include/clang/Serialization/ASTBitCodes.h
clang/lib/AST/ComputeDependence.cpp
clang/lib/AST/Expr.cpp
clang/lib/AST/ExprClassification.cpp
clang/lib/AST/ExprConstant.cpp
clang/lib/AST/ItaniumMangle.cpp
clang/lib/AST/JSONNodeDumper.cpp
clang/lib/AST/OpenACCClause.cpp
clang/lib/AST/StmtPrinter.cpp
clang/lib/AST/StmtProfile.cpp
clang/lib/AST/TextNodeDumper.cpp
clang/lib/CodeGen/CGExprScalar.cpp
clang/lib/Parse/ParseOpenACC.cpp
clang/lib/Sema/SemaExceptionSpec.cpp
clang/lib/Sema/SemaOpenACC.cpp
clang/lib/Sema/TreeTransform.h
clang/lib/Serialization/ASTReader.cpp
clang/lib/Serialization/ASTReaderStmt.cpp
clang/lib/Serialization/ASTWriter.cpp
clang/lib/Serialization/ASTWriterStmt.cpp
clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
clang/test/AST/ast-print-openacc-loop-construct.cpp
clang/test/ParserOpenACC/parse-clauses.c
clang/test/SemaOpenACC/compute-construct-device_type-clause.c
clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
clang/test/SemaOpenACC/loop-construct-device_type-clause.c
clang/tools/libclang/CIndex.cpp
clang/tools/libclang/CXCursor.cpp
Removed:
################################################################################
diff --git a/clang/include/clang/AST/ComputeDependence.h b/clang/include/clang/AST/ComputeDependence.h
index 6d3a51c379f9df..1a8507cfbf9872 100644
--- a/clang/include/clang/AST/ComputeDependence.h
+++ b/clang/include/clang/AST/ComputeDependence.h
@@ -107,6 +107,7 @@ class ObjCSubscriptRefExpr;
class ObjCIsaExpr;
class ObjCIndirectCopyRestoreExpr;
class ObjCMessageExpr;
+class OpenACCAsteriskSizeExpr;
// The following functions are called from constructors of `Expr`, so they
// should not access anything beyond basic
@@ -203,6 +204,7 @@ ExprDependence computeDependence(ObjCSubscriptRefExpr *E);
ExprDependence computeDependence(ObjCIsaExpr *E);
ExprDependence computeDependence(ObjCIndirectCopyRestoreExpr *E);
ExprDependence computeDependence(ObjCMessageExpr *E);
+ExprDependence computeDependence(OpenACCAsteriskSizeExpr *E);
} // namespace clang
#endif
diff --git a/clang/include/clang/AST/Expr.h b/clang/include/clang/AST/Expr.h
index 66c746cc25040f..57353855c51e7c 100644
--- a/clang/include/clang/AST/Expr.h
+++ b/clang/include/clang/AST/Expr.h
@@ -2072,6 +2072,41 @@ class PredefinedExpr final
}
};
+/// This expression type represents an asterisk in an OpenACC Size-Expr, used in
+/// the 'tile' and 'gang' clauses. It is of 'int' type, but should not be
+/// evaluated.
+class OpenACCAsteriskSizeExpr final : public Expr {
+ friend class ASTStmtReader;
+ SourceLocation AsteriskLoc;
+
+ OpenACCAsteriskSizeExpr(SourceLocation AsteriskLoc, QualType IntTy)
+ : Expr(OpenACCAsteriskSizeExprClass, IntTy, VK_PRValue, OK_Ordinary),
+ AsteriskLoc(AsteriskLoc) {}
+
+ void setAsteriskLocation(SourceLocation Loc) { AsteriskLoc = Loc; }
+
+public:
+ static OpenACCAsteriskSizeExpr *Create(const ASTContext &C,
+ SourceLocation Loc);
+ static OpenACCAsteriskSizeExpr *CreateEmpty(const ASTContext &C);
+
+ SourceLocation getBeginLoc() const { return AsteriskLoc; }
+ SourceLocation getEndLoc() const { return AsteriskLoc; }
+ SourceLocation getLocation() const { return AsteriskLoc; }
+
+ static bool classof(const Stmt *T) {
+ return T->getStmtClass() == OpenACCAsteriskSizeExprClass;
+ }
+ // Iterators
+ child_range children() {
+ return child_range(child_iterator(), child_iterator());
+ }
+
+ const_child_range children() const {
+ return const_child_range(const_child_iterator(), const_child_iterator());
+ }
+};
+
// This represents a use of the __builtin_sycl_unique_stable_name, which takes a
// type-id, and at CodeGen time emits a unique string representation of the
// type in a way that permits us to properly encode information about the SYCL
diff --git a/clang/include/clang/AST/JSONNodeDumper.h b/clang/include/clang/AST/JSONNodeDumper.h
index 55bd583e304e8b..9422c8fceccfbd 100644
--- a/clang/include/clang/AST/JSONNodeDumper.h
+++ b/clang/include/clang/AST/JSONNodeDumper.h
@@ -283,6 +283,7 @@ class JSONNodeDumper
void VisitDeclRefExpr(const DeclRefExpr *DRE);
void VisitSYCLUniqueStableNameExpr(const SYCLUniqueStableNameExpr *E);
+ void VisitOpenACCAsteriskSizeExpr(const OpenACCAsteriskSizeExpr *E);
void VisitPredefinedExpr(const PredefinedExpr *PE);
void VisitUnaryOperator(const UnaryOperator *UO);
void VisitBinaryOperator(const BinaryOperator *BO);
diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index 90f5b7fc9ab6f4..e4f2e07222a338 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -481,6 +481,35 @@ class OpenACCNumGangsClause final
}
};
+class OpenACCTileClause final
+ : public OpenACCClauseWithExprs,
+ public llvm::TrailingObjects<OpenACCTileClause, Expr *> {
+ OpenACCTileClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
+ ArrayRef<Expr *> SizeExprs, SourceLocation EndLoc)
+ : OpenACCClauseWithExprs(OpenACCClauseKind::Tile, BeginLoc, LParenLoc,
+ EndLoc) {
+ std::uninitialized_copy(SizeExprs.begin(), SizeExprs.end(),
+ getTrailingObjects<Expr *>());
+ setExprs(MutableArrayRef(getTrailingObjects<Expr *>(), SizeExprs.size()));
+ }
+
+public:
+ static bool classof(const OpenACCClause *C) {
+ return C->getClauseKind() == OpenACCClauseKind::Tile;
+ }
+ static OpenACCTileClause *Create(const ASTContext &C, SourceLocation BeginLoc,
+ SourceLocation LParenLoc,
+ ArrayRef<Expr *> SizeExprs,
+ SourceLocation EndLoc);
+ llvm::ArrayRef<Expr *> getSizeExprs() {
+ return OpenACCClauseWithExprs::getExprs();
+ }
+
+ llvm::ArrayRef<Expr *> getSizeExprs() const {
+ return OpenACCClauseWithExprs::getExprs();
+ }
+};
+
/// Represents one of a handful of clauses that have a single integer
/// expression.
class OpenACCClauseWithSingleIntExpr : public OpenACCClauseWithExprs {
diff --git a/clang/include/clang/AST/RecursiveASTVisitor.h b/clang/include/clang/AST/RecursiveASTVisitor.h
index cd9947f7ab9805..cbbba9e88b7f5b 100644
--- a/clang/include/clang/AST/RecursiveASTVisitor.h
+++ b/clang/include/clang/AST/RecursiveASTVisitor.h
@@ -2867,6 +2867,7 @@ DEF_TRAVERSE_STMT(ParenListExpr, {})
DEF_TRAVERSE_STMT(SYCLUniqueStableNameExpr, {
TRY_TO(TraverseTypeLoc(S->getTypeSourceInfo()->getTypeLoc()));
})
+DEF_TRAVERSE_STMT(OpenACCAsteriskSizeExpr, {})
DEF_TRAVERSE_STMT(PredefinedExpr, {})
DEF_TRAVERSE_STMT(ShuffleVectorExpr, {})
DEF_TRAVERSE_STMT(ConvertVectorExpr, {})
diff --git a/clang/include/clang/AST/TextNodeDumper.h b/clang/include/clang/AST/TextNodeDumper.h
index 57100e7ede171c..9c320c8ae3e54c 100644
--- a/clang/include/clang/AST/TextNodeDumper.h
+++ b/clang/include/clang/AST/TextNodeDumper.h
@@ -410,6 +410,7 @@ class TextNodeDumper
void VisitHLSLOutArgExpr(const HLSLOutArgExpr *E);
void VisitOpenACCConstructStmt(const OpenACCConstructStmt *S);
void VisitOpenACCLoopConstruct(const OpenACCLoopConstruct *S);
+ void VisitOpenACCAsteriskSizeExpr(const OpenACCAsteriskSizeExpr *S);
void VisitEmbedExpr(const EmbedExpr *S);
void VisitAtomicExpr(const AtomicExpr *AE);
};
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index dc84110ef78211..552120f6b7c9d3 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12666,6 +12666,10 @@ def err_acc_loop_spec_conflict
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">;
+def err_acc_size_expr_value
+ : Error<
+ "OpenACC 'tile' clause size expression must be %select{an asterisk "
+ "or a constant expression|positive integer value, evaluated to %1}0">;
def err_acc_invalid_in_collapse_loop
: Error<"%select{OpenACC '%1' construct|while loop|do loop}0 cannot appear "
"in intervening code of a 'loop' with a 'collapse' clause">;
diff --git a/clang/include/clang/Basic/OpenACCClauses.def b/clang/include/clang/Basic/OpenACCClauses.def
index 19cdfe7672133b..a380e5ae69c418 100644
--- a/clang/include/clang/Basic/OpenACCClauses.def
+++ b/clang/include/clang/Basic/OpenACCClauses.def
@@ -52,6 +52,7 @@ VISIT_CLAUSE(Private)
VISIT_CLAUSE(Reduction)
VISIT_CLAUSE(Self)
VISIT_CLAUSE(Seq)
+VISIT_CLAUSE(Tile)
VISIT_CLAUSE(VectorLength)
VISIT_CLAUSE(Wait)
diff --git a/clang/include/clang/Basic/StmtNodes.td b/clang/include/clang/Basic/StmtNodes.td
index 30f2c8f1dbfde8..dc82a4b6d1f777 100644
--- a/clang/include/clang/Basic/StmtNodes.td
+++ b/clang/include/clang/Basic/StmtNodes.td
@@ -308,5 +308,8 @@ def OpenACCAssociatedStmtConstruct
def OpenACCComputeConstruct : StmtNode<OpenACCAssociatedStmtConstruct>;
def OpenACCLoopConstruct : StmtNode<OpenACCAssociatedStmtConstruct>;
+// OpenACC Additional Expressions.
+def OpenACCAsteriskSizeExpr : StmtNode<Expr>;
+
// HLSL Constructs.
def HLSLOutArgExpr : StmtNode<Expr>;
diff --git a/clang/include/clang/Parse/Parser.h b/clang/include/clang/Parse/Parser.h
index eb8a851da7e04e..93e49d395388a6 100644
--- a/clang/include/clang/Parse/Parser.h
+++ b/clang/include/clang/Parse/Parser.h
@@ -3786,10 +3786,13 @@ class Parser : public CodeCompletionHandler {
OpenACCIntExprParseResult ParseOpenACCAsyncArgument(OpenACCDirectiveKind DK,
OpenACCClauseKind CK,
SourceLocation Loc);
+
/// Parses the 'size-expr', which is an integral value, or an asterisk.
- bool ParseOpenACCSizeExpr();
+ /// Asterisk is represented by a OpenACCAsteriskSizeExpr
+ ExprResult ParseOpenACCSizeExpr(OpenACCClauseKind CK);
/// Parses a comma delimited list of 'size-expr's.
- bool ParseOpenACCSizeExprList();
+ bool ParseOpenACCSizeExprList(OpenACCClauseKind CK,
+ llvm::SmallVectorImpl<Expr *> &SizeExprs);
/// Parses a 'gang-arg-list', used for the 'gang' clause.
bool ParseOpenACCGangArgList(SourceLocation GangLoc);
/// Parses a 'gang-arg', used for the 'gang' clause.
diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index 26564835fa1af6..d25c52ec3be43a 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -179,6 +179,7 @@ class SemaOpenACC : public SemaBase {
assert((ClauseKind == OpenACCClauseKind::NumGangs ||
ClauseKind == OpenACCClauseKind::NumWorkers ||
ClauseKind == OpenACCClauseKind::Async ||
+ ClauseKind == OpenACCClauseKind::Tile ||
ClauseKind == OpenACCClauseKind::VectorLength) &&
"Parsed clause kind does not have a int exprs");
@@ -224,6 +225,7 @@ class SemaOpenACC : public SemaBase {
assert((ClauseKind == OpenACCClauseKind::NumGangs ||
ClauseKind == OpenACCClauseKind::NumWorkers ||
ClauseKind == OpenACCClauseKind::Async ||
+ ClauseKind == OpenACCClauseKind::Tile ||
ClauseKind == OpenACCClauseKind::VectorLength) &&
"Parsed clause kind does not have a int exprs");
@@ -335,6 +337,7 @@ class SemaOpenACC : public SemaBase {
assert((ClauseKind == OpenACCClauseKind::NumGangs ||
ClauseKind == OpenACCClauseKind::NumWorkers ||
ClauseKind == OpenACCClauseKind::Async ||
+ ClauseKind == OpenACCClauseKind::Tile ||
ClauseKind == OpenACCClauseKind::VectorLength) &&
"Parsed clause kind does not have a int exprs");
Details = IntExprDetails{{IntExprs.begin(), IntExprs.end()}};
@@ -343,6 +346,7 @@ class SemaOpenACC : public SemaBase {
assert((ClauseKind == OpenACCClauseKind::NumGangs ||
ClauseKind == OpenACCClauseKind::NumWorkers ||
ClauseKind == OpenACCClauseKind::Async ||
+ ClauseKind == OpenACCClauseKind::Tile ||
ClauseKind == OpenACCClauseKind::VectorLength) &&
"Parsed clause kind does not have a int exprs");
Details = IntExprDetails{std::move(IntExprs)};
@@ -522,6 +526,12 @@ class SemaOpenACC : public SemaBase {
SourceLocation RBLoc);
/// Checks the loop depth value for a collapse clause.
ExprResult CheckCollapseLoopCount(Expr *LoopCount);
+ /// Checks a single size expr for a tile clause. 'gang' could possibly call
+ /// this, but has slightly stricter rules as to valid values.
+ ExprResult CheckTileSizeExpr(Expr *SizeExpr);
+
+ ExprResult BuildOpenACCAsteriskSizeExpr(SourceLocation AsteriskLoc);
+ ExprResult ActOnOpenACCAsteriskSizeExpr(SourceLocation AsteriskLoc);
/// Helper type to restore the state of various 'loop' constructs when we run
/// into a loop (for, etc) inside the construct.
diff --git a/clang/include/clang/Serialization/ASTBitCodes.h b/clang/include/clang/Serialization/ASTBitCodes.h
index 1af1f4a10db290..4b79d4b7711905 100644
--- a/clang/include/clang/Serialization/ASTBitCodes.h
+++ b/clang/include/clang/Serialization/ASTBitCodes.h
@@ -2002,12 +2002,14 @@ enum StmtCode {
// SYCLUniqueStableNameExpr
EXPR_SYCL_UNIQUE_STABLE_NAME,
- // OpenACC Constructs
+ // OpenACC Constructs/Exprs
STMT_OPENACC_COMPUTE_CONSTRUCT,
STMT_OPENACC_LOOP_CONSTRUCT,
+ EXPR_OPENACC_ASTERISK_SIZE,
// HLSL Constructs
EXPR_HLSL_OUT_ARG,
+
};
/// The kinds of designators that can occur in a
diff --git a/clang/lib/AST/ComputeDependence.cpp b/clang/lib/AST/ComputeDependence.cpp
index 6ef49790481aca..8c79df8317a2ec 100644
--- a/clang/lib/AST/ComputeDependence.cpp
+++ b/clang/lib/AST/ComputeDependence.cpp
@@ -958,3 +958,9 @@ ExprDependence clang::computeDependence(ObjCMessageExpr *E) {
D |= A->getDependence();
return D;
}
+
+ExprDependence clang::computeDependence(OpenACCAsteriskSizeExpr *E) {
+ // This represents a simple asterisk as typed, so cannot be dependent in any
+ // way.
+ return ExprDependence::None;
+}
diff --git a/clang/lib/AST/Expr.cpp b/clang/lib/AST/Expr.cpp
index 2e463fc00c6b68..9ecbf121e3fc0d 100644
--- a/clang/lib/AST/Expr.cpp
+++ b/clang/lib/AST/Expr.cpp
@@ -3640,6 +3640,7 @@ bool Expr::HasSideEffects(const ASTContext &Ctx,
case SYCLUniqueStableNameExprClass:
case PackIndexingExprClass:
case HLSLOutArgExprClass:
+ case OpenACCAsteriskSizeExprClass:
// These never have a side-effect.
return false;
@@ -5408,3 +5409,13 @@ HLSLOutArgExpr *HLSLOutArgExpr::Create(const ASTContext &C, QualType Ty,
HLSLOutArgExpr *HLSLOutArgExpr::CreateEmpty(const ASTContext &C) {
return new (C) HLSLOutArgExpr(EmptyShell());
}
+
+OpenACCAsteriskSizeExpr *OpenACCAsteriskSizeExpr::Create(const ASTContext &C,
+ SourceLocation Loc) {
+ return new (C) OpenACCAsteriskSizeExpr(Loc, C.IntTy);
+}
+
+OpenACCAsteriskSizeExpr *
+OpenACCAsteriskSizeExpr::CreateEmpty(const ASTContext &C) {
+ return new (C) OpenACCAsteriskSizeExpr({}, C.IntTy);
+}
diff --git a/clang/lib/AST/ExprClassification.cpp b/clang/lib/AST/ExprClassification.cpp
index 9d97633309ada2..3f37d06cc8f3a0 100644
--- a/clang/lib/AST/ExprClassification.cpp
+++ b/clang/lib/AST/ExprClassification.cpp
@@ -471,6 +471,7 @@ static Cl::Kinds ClassifyInternal(ASTContext &Ctx, const Expr *E) {
case Expr::CoyieldExprClass:
return ClassifyInternal(Ctx, cast<CoroutineSuspendExpr>(E)->getResumeExpr());
case Expr::SYCLUniqueStableNameExprClass:
+ case Expr::OpenACCAsteriskSizeExprClass:
return Cl::CL_PRValue;
break;
diff --git a/clang/lib/AST/ExprConstant.cpp b/clang/lib/AST/ExprConstant.cpp
index 3a73cea97fcc32..4d5af96093cfeb 100644
--- a/clang/lib/AST/ExprConstant.cpp
+++ b/clang/lib/AST/ExprConstant.cpp
@@ -11873,6 +11873,13 @@ class IntExprEvaluator
return Success(E->getValue(), E);
}
+ bool VisitOpenACCAsteriskSizeExpr(const OpenACCAsteriskSizeExpr *E) {
+ // This should not be evaluated during constant expr evaluation, as it
+ // should always be in an unevaluated context (the args list of a 'gang' or
+ // 'tile' clause).
+ return Error(E);
+ }
+
bool VisitUnaryReal(const UnaryOperator *E);
bool VisitUnaryImag(const UnaryOperator *E);
@@ -16908,6 +16915,7 @@ static ICEDiag CheckICE(const Expr* E, const ASTContext &Ctx) {
case Expr::GNUNullExprClass:
case Expr::SourceLocExprClass:
case Expr::EmbedExprClass:
+ case Expr::OpenACCAsteriskSizeExprClass:
return NoDiag();
case Expr::PackIndexingExprClass:
diff --git a/clang/lib/AST/ItaniumMangle.cpp b/clang/lib/AST/ItaniumMangle.cpp
index 70ed9cb0736f59..769a863c2b6764 100644
--- a/clang/lib/AST/ItaniumMangle.cpp
+++ b/clang/lib/AST/ItaniumMangle.cpp
@@ -5745,6 +5745,15 @@ void CXXNameMangler::mangleExpression(const Expr *E, unsigned Arity,
case Expr::HLSLOutArgExprClass:
llvm_unreachable(
"cannot mangle hlsl temporary value; mangling wrong thing?");
+ case Expr::OpenACCAsteriskSizeExprClass: {
+ // We shouldn't ever be able to get here, but diagnose anyway.
+ DiagnosticsEngine &Diags = Context.getDiags();
+ unsigned DiagID = Diags.getCustomDiagID(
+ DiagnosticsEngine::Error,
+ "cannot yet mangle OpenACC Asterisk Size expression");
+ Diags.Report(DiagID);
+ return;
+ }
}
if (AsTemplateArg && !IsPrimaryExpr)
diff --git a/clang/lib/AST/JSONNodeDumper.cpp b/clang/lib/AST/JSONNodeDumper.cpp
index 565f1e05710c87..ddbe2136a671f3 100644
--- a/clang/lib/AST/JSONNodeDumper.cpp
+++ b/clang/lib/AST/JSONNodeDumper.cpp
@@ -1354,6 +1354,9 @@ void JSONNodeDumper::VisitSYCLUniqueStableNameExpr(
createQualType(E->getTypeSourceInfo()->getType()));
}
+void JSONNodeDumper::VisitOpenACCAsteriskSizeExpr(
+ const OpenACCAsteriskSizeExpr *E) {}
+
void JSONNodeDumper::VisitPredefinedExpr(const PredefinedExpr *PE) {
JOS.attribute("name", PredefinedExpr::getIdentKindName(PE->getIdentKind()));
}
diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp
index d864ded33e8d1f..0b34ed6189593e 100644
--- a/clang/lib/AST/OpenACCClause.cpp
+++ b/clang/lib/AST/OpenACCClause.cpp
@@ -24,6 +24,7 @@ bool OpenACCClauseWithParams::classof(const OpenACCClause *C) {
}
bool OpenACCClauseWithExprs::classof(const OpenACCClause *C) {
return OpenACCWaitClause::classof(C) || OpenACCNumGangsClause::classof(C) ||
+ OpenACCTileClause::classof(C) ||
OpenACCClauseWithSingleIntExpr::classof(C) ||
OpenACCClauseWithVarList::classof(C);
}
@@ -221,6 +222,16 @@ OpenACCNumGangsClause *OpenACCNumGangsClause::Create(const ASTContext &C,
return new (Mem) OpenACCNumGangsClause(BeginLoc, LParenLoc, IntExprs, EndLoc);
}
+OpenACCTileClause *OpenACCTileClause::Create(const ASTContext &C,
+ SourceLocation BeginLoc,
+ SourceLocation LParenLoc,
+ ArrayRef<Expr *> SizeExprs,
+ SourceLocation EndLoc) {
+ void *Mem =
+ C.Allocate(OpenACCTileClause::totalSizeToAlloc<Expr *>(SizeExprs.size()));
+ return new (Mem) OpenACCTileClause(BeginLoc, LParenLoc, SizeExprs, EndLoc);
+}
+
OpenACCPrivateClause *OpenACCPrivateClause::Create(const ASTContext &C,
SourceLocation BeginLoc,
SourceLocation LParenLoc,
@@ -420,6 +431,13 @@ void OpenACCClausePrinter::VisitNumGangsClause(const OpenACCNumGangsClause &C) {
OS << ")";
}
+void OpenACCClausePrinter::VisitTileClause(const OpenACCTileClause &C) {
+ OS << "tile(";
+ llvm::interleaveComma(C.getSizeExprs(), OS,
+ [&](const Expr *E) { printExpr(E); });
+ OS << ")";
+}
+
void OpenACCClausePrinter::VisitNumWorkersClause(
const OpenACCNumWorkersClause &C) {
OS << "num_workers(";
diff --git a/clang/lib/AST/StmtPrinter.cpp b/clang/lib/AST/StmtPrinter.cpp
index e1b5bec7a50d0a..641f7b52de6dfb 100644
--- a/clang/lib/AST/StmtPrinter.cpp
+++ b/clang/lib/AST/StmtPrinter.cpp
@@ -1309,6 +1309,10 @@ void StmtPrinter::VisitPredefinedExpr(PredefinedExpr *Node) {
OS << PredefinedExpr::getIdentKindName(Node->getIdentKind());
}
+void StmtPrinter::VisitOpenACCAsteriskSizeExpr(OpenACCAsteriskSizeExpr *Node) {
+ OS << '*';
+}
+
void StmtPrinter::VisitCharacterLiteral(CharacterLiteral *Node) {
CharacterLiteral::print(Node->getValue(), Node->getKind(), OS);
}
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index c3812844ab8a31..299ac005c7fdb9 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -1360,6 +1360,11 @@ void StmtProfiler::VisitPredefinedExpr(const PredefinedExpr *S) {
ID.AddInteger(llvm::to_underlying(S->getIdentKind()));
}
+void StmtProfiler::VisitOpenACCAsteriskSizeExpr(
+ const OpenACCAsteriskSizeExpr *S) {
+ VisitExpr(S);
+}
+
void StmtProfiler::VisitIntegerLiteral(const IntegerLiteral *S) {
VisitExpr(S);
S->getValue().Profile(ID);
@@ -2552,6 +2557,11 @@ void OpenACCClauseProfiler::VisitNumGangsClause(
Profiler.VisitStmt(E);
}
+void OpenACCClauseProfiler::VisitTileClause(const OpenACCTileClause &Clause) {
+ for (auto *E : Clause.getSizeExprs())
+ Profiler.VisitStmt(E);
+}
+
void OpenACCClauseProfiler::VisitNumWorkersClause(
const OpenACCNumWorkersClause &Clause) {
assert(Clause.hasIntExpr() && "num_workers clause requires a valid int expr");
diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp
index 8a74159c7c93e5..15b23d60c3ffab 100644
--- a/clang/lib/AST/TextNodeDumper.cpp
+++ b/clang/lib/AST/TextNodeDumper.cpp
@@ -381,6 +381,11 @@ void TextNodeDumper::Visit(const OMPClause *C) {
OS << " <implicit>";
}
+void TextNodeDumper::VisitOpenACCAsteriskSizeExpr(
+ const OpenACCAsteriskSizeExpr *E) {
+ // Nothing to do here, only location exists, and that is printed elsewhere.
+}
+
void TextNodeDumper::Visit(const OpenACCClause *C) {
if (!C) {
ColorScope Color(OS, ShowColors, NullColor);
@@ -414,6 +419,7 @@ void TextNodeDumper::Visit(const OpenACCClause *C) {
case OpenACCClauseKind::Private:
case OpenACCClauseKind::Self:
case OpenACCClauseKind::Seq:
+ case OpenACCClauseKind::Tile:
case OpenACCClauseKind::VectorLength:
// The condition expression will be printed as a part of the 'children',
// but print 'clause' here so it is clear what is happening from the dump.
diff --git a/clang/lib/CodeGen/CGExprScalar.cpp b/clang/lib/CodeGen/CGExprScalar.cpp
index b7f5b932c56b6f..7529d4465f2c34 100644
--- a/clang/lib/CodeGen/CGExprScalar.cpp
+++ b/clang/lib/CodeGen/CGExprScalar.cpp
@@ -531,6 +531,10 @@ class ScalarExprEmitter
return CGF.getOrCreateOpaqueRValueMapping(E).getScalarVal();
}
+ Value *VisitOpenACCAsteriskSizeExpr(OpenACCAsteriskSizeExpr *E) {
+ llvm_unreachable("Codegen for this isn't defined/implemented");
+ }
+
// l-values.
Value *VisitDeclRefExpr(DeclRefExpr *E) {
if (CodeGenFunction::ConstantEmission Constant = CGF.tryEmitAsConstant(E))
diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
index 87673beb34300e..b27e50b147f4a8 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -743,38 +743,50 @@ bool Parser::ParseOpenACCDeviceTypeList(
// int-expr
// Note that this is specified under 'gang-arg-list', but also applies to 'tile'
// via reference.
-bool Parser::ParseOpenACCSizeExpr() {
- // FIXME: Ensure these are constant expressions.
-
+ExprResult Parser::ParseOpenACCSizeExpr(OpenACCClauseKind CK) {
// The size-expr ends up being ambiguous when only looking at the current
// token, as it could be a deref of a variable/expression.
if (getCurToken().is(tok::star) &&
NextToken().isOneOf(tok::comma, tok::r_paren,
tok::annot_pragma_openacc_end)) {
- ConsumeToken();
- return false;
+ SourceLocation AsteriskLoc = ConsumeToken();
+ return getActions().OpenACC().ActOnOpenACCAsteriskSizeExpr(AsteriskLoc);
}
- return getActions()
- .CorrectDelayedTyposInExpr(ParseAssignmentExpression())
- .isInvalid();
+ ExprResult SizeExpr =
+ getActions().CorrectDelayedTyposInExpr(ParseConstantExpression());
+
+ if (!SizeExpr.isUsable())
+ return SizeExpr;
+
+ SizeExpr = getActions().OpenACC().ActOnIntExpr(
+ OpenACCDirectiveKind::Invalid, CK, SizeExpr.get()->getBeginLoc(),
+ SizeExpr.get());
+
+ return SizeExpr;
}
-bool Parser::ParseOpenACCSizeExprList() {
- if (ParseOpenACCSizeExpr()) {
+bool Parser::ParseOpenACCSizeExprList(
+ OpenACCClauseKind CK, llvm::SmallVectorImpl<Expr *> &SizeExprs) {
+ ExprResult SizeExpr = ParseOpenACCSizeExpr(CK);
+ if (!SizeExpr.isUsable()) {
SkipUntil(tok::r_paren, tok::annot_pragma_openacc_end,
Parser::StopBeforeMatch);
- return false;
+ return true;
}
+ SizeExprs.push_back(SizeExpr.get());
+
while (!getCurToken().isOneOf(tok::r_paren, tok::annot_pragma_openacc_end)) {
ExpectAndConsume(tok::comma);
- if (ParseOpenACCSizeExpr()) {
+ SizeExpr = ParseOpenACCSizeExpr(CK);
+ if (!SizeExpr.isUsable()) {
SkipUntil(tok::r_paren, tok::annot_pragma_openacc_end,
Parser::StopBeforeMatch);
- return false;
+ return true;
}
+ SizeExprs.push_back(SizeExpr.get());
}
return false;
}
@@ -792,7 +804,7 @@ bool Parser::ParseOpenACCGangArg(SourceLocation GangLoc) {
// 'static' just takes a size-expr, which is an int-expr or an asterisk.
ConsumeToken();
ConsumeToken();
- return ParseOpenACCSizeExpr();
+ return ParseOpenACCSizeExpr(OpenACCClauseKind::Gang).isInvalid();
}
if (isOpenACCSpecialToken(OpenACCSpecialTokenKind::Dim, getCurToken()) &&
@@ -1052,12 +1064,16 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
}
break;
}
- case OpenACCClauseKind::Tile:
- if (ParseOpenACCSizeExprList()) {
+ case OpenACCClauseKind::Tile: {
+ llvm::SmallVector<Expr *> SizeExprs;
+ if (ParseOpenACCSizeExprList(OpenACCClauseKind::Tile, SizeExprs)) {
Parens.skipToEnd();
return OpenACCCanContinue();
}
+
+ ParsedClause.setIntExprDetails(std::move(SizeExprs));
break;
+ }
default:
llvm_unreachable("Not a required parens type?");
}
diff --git a/clang/lib/Sema/SemaExceptionSpec.cpp b/clang/lib/Sema/SemaExceptionSpec.cpp
index 8aedbfcf878a11..dbddd6c370aa07 100644
--- a/clang/lib/Sema/SemaExceptionSpec.cpp
+++ b/clang/lib/Sema/SemaExceptionSpec.cpp
@@ -1593,6 +1593,8 @@ CanThrowResult Sema::canThrow(const Stmt *S) {
case Stmt::SYCLUniqueStableNameExprClass:
return CT_Cannot;
+ case Stmt::OpenACCAsteriskSizeExprClass:
+ return CT_Cannot;
case Stmt::NoStmtClass:
llvm_unreachable("Invalid class for statement");
}
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index dfaf726891cfc8..f3840c6ffba82b 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -354,6 +354,17 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
return false;
}
}
+ case OpenACCClauseKind::Tile: {
+ 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.
@@ -526,6 +537,36 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitDefaultClause(
Clause.getLParenLoc(), Clause.getEndLoc());
}
+OpenACCClause *SemaOpenACCClauseVisitor::VisitTileClause(
+ SemaOpenACC::OpenACCParsedClause &Clause) {
+ if (Clause.getDirectiveKind() != OpenACCDirectiveKind::Loop)
+ return isNotImplemented();
+
+ // 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;
+
+ llvm::SmallVector<Expr *> NewSizeExprs;
+
+ // Make sure these are all positive constant expressions or *.
+ for (Expr *E : Clause.getIntExprs()) {
+ ExprResult Res = SemaRef.CheckTileSizeExpr(E);
+
+ if (!Res.isUsable())
+ return nullptr;
+
+ NewSizeExprs.push_back(Res.get());
+ }
+
+ return OpenACCTileClause::Create(Ctx, Clause.getBeginLoc(),
+ Clause.getLParenLoc(), NewSizeExprs,
+ Clause.getEndLoc());
+}
+
OpenACCClause *SemaOpenACCClauseVisitor::VisitIfClause(
SemaOpenACC::OpenACCParsedClause &Clause) {
// Restrictions only properly implemented on 'compute' constructs, and
@@ -1698,6 +1739,34 @@ ExprResult SemaOpenACC::CheckCollapseLoopCount(Expr *LoopCount) {
ConstantExpr::Create(getASTContext(), LoopCount, APValue{*ICE})};
}
+ExprResult SemaOpenACC::CheckTileSizeExpr(Expr *SizeExpr) {
+ if (!SizeExpr)
+ return ExprError();
+
+ assert((SizeExpr->isInstantiationDependent() ||
+ SizeExpr->getType()->isIntegerType()) &&
+ "size argument non integer?");
+
+ // If dependent, or an asterisk, the expression is fine.
+ if (SizeExpr->isInstantiationDependent() ||
+ isa<OpenACCAsteriskSizeExpr>(SizeExpr))
+ return ExprResult{SizeExpr};
+
+ std::optional<llvm::APSInt> ICE =
+ SizeExpr->getIntegerConstantExpr(getASTContext());
+
+ // OpenACC 3.3 2.9.8
+ // where each tile size is a constant positive integer expression or asterisk.
+ if (!ICE || *ICE <= 0) {
+ Diag(SizeExpr->getBeginLoc(), diag::err_acc_size_expr_value)
+ << ICE.has_value() << ICE.value_or(llvm::APSInt{}).getExtValue();
+ return ExprError();
+ }
+
+ return ExprResult{
+ ConstantExpr::Create(getASTContext(), SizeExpr, APValue{*ICE})};
+}
+
void SemaOpenACC::ActOnWhileStmt(SourceLocation WhileLoc) {
if (!getLangOpts().OpenACC)
return;
@@ -1944,3 +2013,13 @@ bool SemaOpenACC::ActOnStartDeclDirective(OpenACCDirectiveKind K,
}
DeclGroupRef SemaOpenACC::ActOnEndDeclDirective() { return DeclGroupRef{}; }
+
+ExprResult
+SemaOpenACC::BuildOpenACCAsteriskSizeExpr(SourceLocation AsteriskLoc) {
+ return OpenACCAsteriskSizeExpr::Create(getASTContext(), AsteriskLoc);
+}
+
+ExprResult
+SemaOpenACC::ActOnOpenACCAsteriskSizeExpr(SourceLocation AsteriskLoc) {
+ return BuildOpenACCAsteriskSizeExpr(AsteriskLoc);
+}
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index 04308b2b57d8a0..76efc9fe831854 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -4090,6 +4090,10 @@ class TreeTransform {
OpenACCDirectiveKind::Loop, BeginLoc, DirLoc, EndLoc, Clauses, Loop);
}
+ ExprResult RebuildOpenACCAsteriskSizeExpr(SourceLocation AsteriskLoc) {
+ return getSema().OpenACC().ActOnOpenACCAsteriskSizeExpr(AsteriskLoc);
+ }
+
private:
TypeLoc TransformTypeInObjectScope(TypeLoc TL,
QualType ObjectType,
@@ -11871,6 +11875,36 @@ void OpenACCClauseTransform<Derived>::VisitCollapseClause(
ParsedClause.getLParenLoc(), ParsedClause.isForce(),
ParsedClause.getLoopCount(), ParsedClause.getEndLoc());
}
+
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitTileClause(
+ const OpenACCTileClause &C) {
+
+ llvm::SmallVector<Expr *> TransformedExprs;
+
+ for (Expr *E : C.getSizeExprs()) {
+ ExprResult NewSizeExpr = Self.TransformExpr(E);
+
+ if (!NewSizeExpr.isUsable())
+ return;
+
+ NewSizeExpr = Self.getSema().OpenACC().ActOnIntExpr(
+ OpenACCDirectiveKind::Invalid, ParsedClause.getClauseKind(),
+ NewSizeExpr.get()->getBeginLoc(), NewSizeExpr.get());
+
+ NewSizeExpr = Self.getSema().OpenACC().CheckTileSizeExpr(NewSizeExpr.get());
+
+ if (!NewSizeExpr.isUsable())
+ return;
+ TransformedExprs.push_back(NewSizeExpr.get());
+ }
+
+ ParsedClause.setIntExprDetails(TransformedExprs);
+ NewClause = OpenACCTileClause::Create(
+ Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
+ ParsedClause.getLParenLoc(), ParsedClause.getIntExprs(),
+ ParsedClause.getEndLoc());
+}
} // namespace
template <typename Derived>
OpenACCClause *TreeTransform<Derived>::TransformOpenACCClause(
@@ -11957,6 +11991,15 @@ TreeTransform<Derived>::TransformOpenACCLoopConstruct(OpenACCLoopConstruct *C) {
TransformedClauses, Loop);
}
+template <typename Derived>
+ExprResult TreeTransform<Derived>::TransformOpenACCAsteriskSizeExpr(
+ OpenACCAsteriskSizeExpr *E) {
+ if (getDerived().AlwaysRebuild())
+ return getDerived().RebuildOpenACCAsteriskSizeExpr(E->getLocation());
+ // Nothing can ever change, so there is never anything to transform.
+ return E;
+}
+
//===----------------------------------------------------------------------===//
// Expression transformation
//===----------------------------------------------------------------------===//
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index 4b599b7dbf3ce8..97b79bd1381c02 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -12306,6 +12306,15 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
return OpenACCCollapseClause::Create(getContext(), BeginLoc, LParenLoc,
HasForce, LoopCount, EndLoc);
}
+ case OpenACCClauseKind::Tile: {
+ SourceLocation LParenLoc = readSourceLocation();
+ unsigned NumClauses = readInt();
+ llvm::SmallVector<Expr *> SizeExprs;
+ for (unsigned I = 0; I < NumClauses; ++I)
+ SizeExprs.push_back(readSubExpr());
+ return OpenACCTileClause::Create(getContext(), BeginLoc, LParenLoc,
+ SizeExprs, EndLoc);
+ }
case OpenACCClauseKind::Finalize:
case OpenACCClauseKind::IfPresent:
@@ -12322,7 +12331,6 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
case OpenACCClauseKind::Bind:
case OpenACCClauseKind::DeviceNum:
case OpenACCClauseKind::DefaultAsync:
- case OpenACCClauseKind::Tile:
case OpenACCClauseKind::Gang:
case OpenACCClauseKind::Invalid:
llvm_unreachable("Clause serialization not yet implemented");
diff --git a/clang/lib/Serialization/ASTReaderStmt.cpp b/clang/lib/Serialization/ASTReaderStmt.cpp
index 84743a52d4c8b8..2038fe7829af9b 100644
--- a/clang/lib/Serialization/ASTReaderStmt.cpp
+++ b/clang/lib/Serialization/ASTReaderStmt.cpp
@@ -575,6 +575,11 @@ void ASTStmtReader::VisitConstantExpr(ConstantExpr *E) {
E->setSubExpr(Record.readSubExpr());
}
+void ASTStmtReader::VisitOpenACCAsteriskSizeExpr(OpenACCAsteriskSizeExpr *E) {
+ VisitExpr(E);
+ E->setAsteriskLocation(readSourceLocation());
+}
+
void ASTStmtReader::VisitSYCLUniqueStableNameExpr(SYCLUniqueStableNameExpr *E) {
VisitExpr(E);
@@ -3054,6 +3059,10 @@ Stmt *ASTReader::ReadStmtFromStream(ModuleFile &F) {
S = SYCLUniqueStableNameExpr::CreateEmpty(Context);
break;
+ case EXPR_OPENACC_ASTERISK_SIZE:
+ S = OpenACCAsteriskSizeExpr::CreateEmpty(Context);
+ break;
+
case EXPR_PREDEFINED:
S = PredefinedExpr::CreateEmpty(
Context,
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index 375ddc90482b27..836532ca402ffc 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -8167,6 +8167,14 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
AddStmt(const_cast<Expr *>(CC->getLoopCount()));
return;
}
+ case OpenACCClauseKind::Tile: {
+ const auto *TC = cast<OpenACCTileClause>(C);
+ writeSourceLocation(TC->getLParenLoc());
+ writeUInt32(TC->getSizeExprs().size());
+ for (Expr *E : TC->getSizeExprs())
+ AddStmt(E);
+ return;
+ }
case OpenACCClauseKind::Finalize:
case OpenACCClauseKind::IfPresent:
@@ -8183,7 +8191,6 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
case OpenACCClauseKind::Bind:
case OpenACCClauseKind::DeviceNum:
case OpenACCClauseKind::DefaultAsync:
- case OpenACCClauseKind::Tile:
case OpenACCClauseKind::Gang:
case OpenACCClauseKind::Invalid:
llvm_unreachable("Clause serialization not yet implemented");
diff --git a/clang/lib/Serialization/ASTWriterStmt.cpp b/clang/lib/Serialization/ASTWriterStmt.cpp
index 837136600181c1..64e4894dc29fb7 100644
--- a/clang/lib/Serialization/ASTWriterStmt.cpp
+++ b/clang/lib/Serialization/ASTWriterStmt.cpp
@@ -643,6 +643,12 @@ void ASTStmtWriter::VisitConstantExpr(ConstantExpr *E) {
Code = serialization::EXPR_CONSTANT;
}
+void ASTStmtWriter::VisitOpenACCAsteriskSizeExpr(OpenACCAsteriskSizeExpr *E) {
+ VisitExpr(E);
+ Record.AddSourceLocation(E->getLocation());
+ Code = serialization::EXPR_OPENACC_ASTERISK_SIZE;
+}
+
void ASTStmtWriter::VisitSYCLUniqueStableNameExpr(SYCLUniqueStableNameExpr *E) {
VisitExpr(E);
diff --git a/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp b/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
index fdabba46992b08..b1919d7027cf4d 100644
--- a/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
+++ b/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
@@ -1966,6 +1966,7 @@ void ExprEngine::Visit(const Stmt *S, ExplodedNode *Pred,
case Stmt::OMPArrayShapingExprClass:
case Stmt::OMPIteratorExprClass:
case Stmt::SYCLUniqueStableNameExprClass:
+ case Stmt::OpenACCAsteriskSizeExprClass:
case Stmt::TypeTraitExprClass: {
Bldr.takeNodes(Pred);
ExplodedNodeSet preVisit;
diff --git a/clang/test/AST/ast-print-openacc-loop-construct.cpp b/clang/test/AST/ast-print-openacc-loop-construct.cpp
index ae1f7964f019eb..aee4591cab428f 100644
--- a/clang/test/AST/ast-print-openacc-loop-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-loop-construct.cpp
@@ -2,6 +2,7 @@
struct SomeStruct{};
+constexpr int get_value() { return 1; }
void foo() {
// CHECK: #pragma acc loop
// CHECK-NEXT: for (;;)
@@ -82,4 +83,16 @@ void foo() {
#pragma acc loop collapse(force:2)
for(;;)
for(;;);
+
+// CHECK: #pragma acc loop tile(1, 3, *, get_value())
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc loop tile(1, 3, *, get_value())
+ for(;;)
+ for(;;)
+ for(;;)
+ for(;;);
}
diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c
index 6c9ce4ad5e1969..a752b2ebd18b6a 100644
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@@ -1153,8 +1153,7 @@ void Tile() {
// expected-note at +1{{to match this '('}}
#pragma acc loop tile(
for(;;){}
- // expected-error at +2{{expected expression}}
- // expected-warning at +1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
+ // expected-error at +1{{expected expression}}
#pragma acc loop tile()
for(;;){}
// expected-error at +3{{expected expression}}
@@ -1162,41 +1161,32 @@ void Tile() {
// expected-note at +1{{to match this '('}}
#pragma acc loop tile(,
for(;;){}
- // expected-error at +2{{expected expression}}
- // expected-warning at +1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
+ // expected-error at +1{{expected expression}}
#pragma acc loop tile(,)
for(;;){}
- // expected-error at +2{{use of undeclared identifier 'invalid'}}
- // expected-warning at +1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
+ // expected-error at +1{{use of undeclared identifier 'invalid'}}
#pragma acc loop tile(returns_int(), *, invalid, *)
for(;;){}
- // expected-error at +2{{expected expression}}
- // expected-warning at +1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
+ // expected-error at +1{{expected expression}}
#pragma acc loop tile(returns_int() *, Foo, *)
for(;;){}
- // expected-error at +2{{indirection requires pointer operand ('int' invalid)}}
- // expected-warning at +1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
+ // expected-error at +1{{indirection requires pointer operand ('int' invalid)}}
#pragma acc loop tile(* returns_int() , *)
for(;;){}
- // expected-warning at +1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
#pragma acc loop tile(*)
for(;;){}
- // expected-warning at +1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
+ // expected-error at +1{{OpenACC 'tile' clause size expression must be an asterisk or a constant expression}}
#pragma acc loop tile(*Foo, *Foo)
for(;;){}
- // expected-warning at +1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
#pragma acc loop tile(5)
for(;;){}
- // expected-warning at +1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
#pragma acc loop tile(*, 5)
for(;;){}
- // expected-warning at +1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
#pragma acc loop tile(5, *)
for(;;){}
- // expected-warning at +1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
#pragma acc loop tile(5, *, 3, *)
for(;;){}
}
diff --git a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
index 26f0315fb86f1a..d08497a7782edb 100644
--- a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
@@ -211,8 +211,7 @@ void uses() {
while(1);
#pragma acc kernels device_type(*) async
while(1);
- // expected-error at +2{{OpenACC clause 'tile' may not follow a 'device_type' clause in a compute construct}}
- // expected-note at +1{{previous clause is here}}
+ // expected-error at +1{{OpenACC 'tile' clause is not valid on 'kernels' directive}}
#pragma acc kernels device_type(*) tile(Var, 1)
while(1);
// expected-error at +2{{OpenACC clause 'gang' may not follow a 'dtype' 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 3212c19d089fc9..f66ce6991acb1a 100644
--- a/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
+++ b/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
@@ -165,8 +165,7 @@ void uses() {
// expected-error at +1{{OpenACC 'async' clause is not valid on 'loop' directive}}
#pragma acc loop auto async
for(;;);
- // expected-warning at +1{{OpenACC clause 'tile' not yet implemented}}
-#pragma acc loop auto tile(Var, 1)
+#pragma acc loop auto tile(1+2, 1)
for(;;);
// expected-warning at +1{{OpenACC clause 'gang' not yet implemented}}
#pragma acc loop auto gang
@@ -303,8 +302,7 @@ void uses() {
// expected-error at +1{{OpenACC 'async' clause is not valid on 'loop' directive}}
#pragma acc loop async auto
for(;;);
- // expected-warning at +1{{OpenACC clause 'tile' not yet implemented}}
-#pragma acc loop tile(Var, 1) auto
+#pragma acc loop tile(1+2, 1) auto
for(;;);
// expected-warning at +1{{OpenACC clause 'gang' not yet implemented}}
#pragma acc loop gang auto
@@ -442,8 +440,7 @@ void uses() {
// expected-error at +1{{OpenACC 'async' clause is not valid on 'loop' directive}}
#pragma acc loop independent async
for(;;);
- // expected-warning at +1{{OpenACC clause 'tile' not yet implemented}}
-#pragma acc loop independent tile(Var, 1)
+#pragma acc loop independent tile(1+2, 1)
for(;;);
// expected-warning at +1{{OpenACC clause 'gang' not yet implemented}}
#pragma acc loop independent gang
@@ -580,8 +577,7 @@ void uses() {
// expected-error at +1{{OpenACC 'async' clause is not valid on 'loop' directive}}
#pragma acc loop async independent
for(;;);
- // expected-warning at +1{{OpenACC clause 'tile' not yet implemented}}
-#pragma acc loop tile(Var, 1) independent
+#pragma acc loop tile(1+2, 1) independent
for(;;);
// expected-warning at +1{{OpenACC clause 'gang' not yet implemented}}
#pragma acc loop gang independent
@@ -728,8 +724,7 @@ void uses() {
// expected-error at +1{{OpenACC 'async' clause is not valid on 'loop' directive}}
#pragma acc loop seq async
for(;;);
- // expected-warning at +1{{OpenACC clause 'tile' not yet implemented}}
-#pragma acc loop seq tile(Var, 1)
+#pragma acc loop seq tile(1+2, 1)
for(;;);
// expected-error at +1{{OpenACC 'wait' clause is not valid on 'loop' directive}}
#pragma acc loop seq wait
@@ -875,8 +870,7 @@ void uses() {
// expected-error at +1{{OpenACC 'async' clause is not valid on 'loop' directive}}
#pragma acc loop async seq
for(;;);
- // expected-warning at +1{{OpenACC clause 'tile' not yet implemented}}
-#pragma acc loop tile(Var, 1) seq
+#pragma acc loop tile(1+2, 1) seq
for(;;);
// expected-error at +1{{OpenACC 'wait' clause is not valid on 'loop' directive}}
#pragma acc loop wait seq
diff --git a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
index 47c9239f4f0e96..aa7c3201c8319c 100644
--- a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
@@ -188,8 +188,8 @@ void uses() {
// expected-error at +1{{OpenACC 'async' clause is not valid on 'loop' directive}}
#pragma acc loop device_type(*) async
for(;;);
- // expected-warning at +1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
-#pragma acc loop device_type(*) tile(Var, 1)
+
+#pragma acc loop device_type(*) tile(*, 1)
for(;;);
// expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
#pragma acc loop dtype(*) gang
diff --git a/clang/test/SemaOpenACC/loop-construct-tile-ast.cpp b/clang/test/SemaOpenACC/loop-construct-tile-ast.cpp
new file mode 100644
index 00000000000000..977c7e3b6d4cf5
--- /dev/null
+++ b/clang/test/SemaOpenACC/loop-construct-tile-ast.cpp
@@ -0,0 +1,231 @@
+// 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 tile(S{}, 1, 2, *)
+ for(;;)
+ for(;;)
+ for(;;)
+ for(;;);
+ // CHECK-NEXT: OpenACCLoopConstruct
+ // CHECK-NEXT: tile clause
+ // CHECK-NEXT: ConstantExpr{{.*}} 'int'
+ // CHECK-NEXT: value: Int 1
+ // 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: ConstantExpr{{.*}} 'int'
+ // CHECK-NEXT: value: Int 1
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
+ // CHECK-NEXT: ConstantExpr{{.*}} 'int'
+ // CHECK-NEXT: value: Int 2
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 2
+ // CHECK-NEXT: OpenACCAsteriskSizeExpr{{.*}}'int'
+ // 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: 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
+}
+
+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 tile(S{}, T{}, *, T{} + S{}, Value, Value + 3)
+ for(;;)
+ for(;;)
+ for(;;)
+ for(;;)
+ for(;;)
+ for(;;);
+ // CHECK-NEXT: OpenACCLoopConstruct
+ // CHECK-NEXT: tile clause
+ //
+ // CHECK-NEXT: ConstantExpr{{.*}} 'int'
+ // CHECK-NEXT: value: Int 1
+ // 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: CXXUnresolvedConstructExpr{{.*}} 'T' 'T' list
+ // CHECK-NEXT: InitListExpr{{.*}}'void'
+ //
+ // CHECK-NEXT: OpenACCAsteriskSizeExpr{{.*}}'int'
+ //
+ // CHECK-NEXT: BinaryOperator{{.*}}'<dependent type>' '+'
+ // CHECK-NEXT: CXXUnresolvedConstructExpr{{.*}} 'T' 'T' list
+ // CHECK-NEXT: InitListExpr{{.*}}'void'
+ // CHECK-NEXT: CXXTemporaryObjectExpr{{.*}} 'S' 'void ()' list
+ //
+ // CHECK-NEXT: DeclRefExpr{{.*}}'unsigned int' NonTypeTemplateParm{{.*}}'Value' 'unsigned int'
+ //
+ // CHECK-NEXT: BinaryOperator{{.*}}'unsigned int' '+'
+ // CHECK-NEXT: DeclRefExpr{{.*}}'unsigned int' NonTypeTemplateParm{{.*}}'Value' 'unsigned int'
+ // CHECK-NEXT: ImplicitCastExpr{{.*}} 'unsigned int' <IntegralCast>
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 3
+ // 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: 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: 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: tile clause
+ //
+ // CHECK-NEXT: ConstantExpr{{.*}} 'int'
+ // CHECK-NEXT: value: Int 1
+ // 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: ConstantExpr{{.*}} 'int'
+ // CHECK-NEXT: value: Int 1
+ // 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: OpenACCAsteriskSizeExpr{{.*}}'int'
+ //
+ // CHECK-NEXT: ConstantExpr{{.*}} 'int'
+ // CHECK-NEXT: value: Int 2
+ // CHECK-NEXT: BinaryOperator{{.*}}'int' '+'
+ // 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: ConstantExpr{{.*}} 'unsigned int'
+ // CHECK-NEXT: value: Int 2
+ // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}} 'unsigned int'
+ // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'unsigned int' depth 0 index 1 Value
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'unsigned int' 2
+ //
+ // CHECK-NEXT: ConstantExpr{{.*}} 'unsigned int'
+ // CHECK-NEXT: value: Int 5
+ // CHECK-NEXT: BinaryOperator{{.*}}'unsigned int' '+'
+ // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}} 'unsigned int'
+ // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'unsigned int' depth 0 index 1 Value
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'unsigned int' 2
+ // CHECK-NEXT: ImplicitCastExpr{{.*}} 'unsigned int' <IntegralCast>
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 3
+
+ // 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: 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: 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-tile-clause.cpp b/clang/test/SemaOpenACC/loop-construct-tile-clause.cpp
new file mode 100644
index 00000000000000..28eadde7416c35
--- /dev/null
+++ b/clang/test/SemaOpenACC/loop-construct-tile-clause.cpp
@@ -0,0 +1,95 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+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_zero_constexpr_templ() {
+ // expected-error at +1 2{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to 0}}
+#pragma acc loop tile(*, T{})
+ for(;;)
+ for(;;);
+
+ // expected-error at +1{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to -1}}
+#pragma acc loop tile(Val, *)
+ for(;;)
+ for(;;);
+
+ // expected-error at +1{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to 0}}
+#pragma acc loop tile(zero(), *)
+ for(;;)
+ for(;;);
+}
+
+void negative_zero_constexpr() {
+ negative_zero_constexpr_templ<int, 1>(); // expected-note{{in instantiation of function template specialization}}
+ negative_zero_constexpr_templ<int, -1>(); // expected-note{{in instantiation of function template specialization}}
+
+ // expected-error at +1{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to 0}}
+#pragma acc loop tile(0, *)
+ for(;;)
+ for(;;);
+
+ // expected-error at +1{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to 0}}
+#pragma acc loop tile(1, 0)
+ for(;;)
+ for(;;);
+
+ // expected-error at +1{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to -1}}
+#pragma acc loop tile(1, -1)
+ for(;;)
+ for(;;);
+
+ // expected-error at +1{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to -1}}
+#pragma acc loop tile(-1, 0)
+ for(;;)
+ for(;;);
+
+ // expected-error at +1{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to 0}}
+#pragma acc loop tile(zero(), 0)
+ for(;;)
+ for(;;);
+
+ // expected-error at +1{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to -1}}
+#pragma acc loop tile(1, neg())
+ for(;;)
+ for(;;);
+
+ // expected-error at +1{{OpenACC 'tile' clause size expression must be an asterisk or a constant expression}}
+#pragma acc loop tile(NotConstexpr{})
+ for(;;);
+
+ // expected-error at +1{{OpenACC 'tile' clause size expression must be positive integer value, evaluated to -1}}
+#pragma acc loop tile(1, ConvertsNegative{})
+ for(;;)
+ for(;;);
+
+#pragma acc loop tile(*, ConvertsOne{})
+ for(;;)
+ for(;;);
+
+}
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index d188f794bad20e..e6c323775c999e 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2839,6 +2839,11 @@ void OpenACCClauseEnqueue::VisitNumGangsClause(const OpenACCNumGangsClause &C) {
Visitor.AddStmt(IE);
}
+void OpenACCClauseEnqueue::VisitTileClause(const OpenACCTileClause &C) {
+ for (Expr *IE : C.getSizeExprs())
+ Visitor.AddStmt(IE);
+}
+
void OpenACCClauseEnqueue::VisitPrivateClause(const OpenACCPrivateClause &C) {
VisitVarList(C);
}
diff --git a/clang/tools/libclang/CXCursor.cpp b/clang/tools/libclang/CXCursor.cpp
index 4e068f272a153f..562228cc334f3a 100644
--- a/clang/tools/libclang/CXCursor.cpp
+++ b/clang/tools/libclang/CXCursor.cpp
@@ -337,6 +337,7 @@ CXCursor cxcursor::MakeCXCursor(const Stmt *S, const Decl *Parent,
case Stmt::SYCLUniqueStableNameExprClass:
case Stmt::EmbedExprClass:
case Stmt::HLSLOutArgExprClass:
+ case Stmt::OpenACCAsteriskSizeExprClass:
K = CXCursor_UnexposedExpr;
break;
More information about the cfe-commits
mailing list