[llvm] 6bd8dc9 - [OpenMP] Initial parsing/sema for the 'omp target teams loop' construct
Mike Rice via llvm-commits
llvm-commits at lists.llvm.org
Fri Mar 18 14:06:34 PDT 2022
Author: Mike Rice
Date: 2022-03-18T13:48:32-07:00
New Revision: 6bd8dc91b89d0b130d8c03174b7b3b0d1cf3b331
URL: https://github.com/llvm/llvm-project/commit/6bd8dc91b89d0b130d8c03174b7b3b0d1cf3b331
DIFF: https://github.com/llvm/llvm-project/commit/6bd8dc91b89d0b130d8c03174b7b3b0d1cf3b331.diff
LOG: [OpenMP] Initial parsing/sema for the 'omp target teams loop' construct
Adds basic parsing/sema/serialization support for the
#pragma omp target teams loop directive.
Differential Revision: https://reviews.llvm.org/D122028
Added:
clang/test/OpenMP/target_teams_generic_loop_ast_print.cpp
clang/test/OpenMP/target_teams_generic_loop_messages.cpp
Modified:
clang/include/clang-c/Index.h
clang/include/clang/AST/RecursiveASTVisitor.h
clang/include/clang/AST/StmtOpenMP.h
clang/include/clang/Basic/StmtNodes.td
clang/include/clang/Sema/Sema.h
clang/include/clang/Serialization/ASTBitCodes.h
clang/lib/AST/StmtOpenMP.cpp
clang/lib/AST/StmtPrinter.cpp
clang/lib/AST/StmtProfile.cpp
clang/lib/Basic/OpenMPKinds.cpp
clang/lib/CodeGen/CGStmt.cpp
clang/lib/Parse/ParseOpenMP.cpp
clang/lib/Sema/SemaExceptionSpec.cpp
clang/lib/Sema/SemaOpenMP.cpp
clang/lib/Sema/TreeTransform.h
clang/lib/Serialization/ASTReaderStmt.cpp
clang/lib/Serialization/ASTWriterStmt.cpp
clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
clang/test/Analysis/cfg-openmp.cpp
clang/tools/libclang/CIndex.cpp
clang/tools/libclang/CXCursor.cpp
llvm/include/llvm/Frontend/OpenMP/OMP.td
Removed:
################################################################################
diff --git a/clang/include/clang-c/Index.h b/clang/include/clang-c/Index.h
index 4f4bf8fe14ea0..23ecac4895bd9 100644
--- a/clang/include/clang-c/Index.h
+++ b/clang/include/clang-c/Index.h
@@ -2604,7 +2604,11 @@ enum CXCursorKind {
*/
CXCursor_OMPTeamsGenericLoopDirective = 296,
- CXCursor_LastStmt = CXCursor_OMPTeamsGenericLoopDirective,
+ /** OpenMP target teams loop directive.
+ */
+ CXCursor_OMPTargetTeamsGenericLoopDirective = 297,
+
+ CXCursor_LastStmt = CXCursor_OMPTargetTeamsGenericLoopDirective,
/**
* Cursor that represents the translation unit itself.
diff --git a/clang/include/clang/AST/RecursiveASTVisitor.h b/clang/include/clang/AST/RecursiveASTVisitor.h
index ef2fb6cbe03ee..b614b6a8bebe0 100644
--- a/clang/include/clang/AST/RecursiveASTVisitor.h
+++ b/clang/include/clang/AST/RecursiveASTVisitor.h
@@ -3077,6 +3077,9 @@ DEF_TRAVERSE_STMT(OMPGenericLoopDirective,
DEF_TRAVERSE_STMT(OMPTeamsGenericLoopDirective,
{ TRY_TO(TraverseOMPExecutableDirective(S)); })
+DEF_TRAVERSE_STMT(OMPTargetTeamsGenericLoopDirective,
+ { TRY_TO(TraverseOMPExecutableDirective(S)); })
+
// OpenMP clauses.
template <typename Derived>
bool RecursiveASTVisitor<Derived>::TraverseOMPClause(OMPClause *C) {
diff --git a/clang/include/clang/AST/StmtOpenMP.h b/clang/include/clang/AST/StmtOpenMP.h
index 74c8d6d980c1e..82d9a69035abd 100644
--- a/clang/include/clang/AST/StmtOpenMP.h
+++ b/clang/include/clang/AST/StmtOpenMP.h
@@ -1527,6 +1527,7 @@ class OMPLoopDirective : public OMPLoopBasedDirective {
T->getStmtClass() == OMPMasterTaskLoopSimdDirectiveClass ||
T->getStmtClass() == OMPGenericLoopDirectiveClass ||
T->getStmtClass() == OMPTeamsGenericLoopDirectiveClass ||
+ T->getStmtClass() == OMPTargetTeamsGenericLoopDirectiveClass ||
T->getStmtClass() == OMPParallelMasterTaskLoopDirectiveClass ||
T->getStmtClass() == OMPParallelMasterTaskLoopSimdDirectiveClass ||
T->getStmtClass() == OMPDistributeDirectiveClass ||
@@ -5635,6 +5636,71 @@ class OMPTeamsGenericLoopDirective final : public OMPLoopDirective {
}
};
+/// This represents '#pragma omp target teams loop' directive.
+///
+/// \code
+/// #pragma omp target teams loop private(a,b) order(concurrent)
+/// \endcode
+/// In this example directive '#pragma omp target teams loop' has
+/// clauses 'private' with the variables 'a' and 'b', and order(concurrent).
+///
+class OMPTargetTeamsGenericLoopDirective final : public OMPLoopDirective {
+ friend class ASTStmtReader;
+ friend class OMPExecutableDirective;
+ /// Build directive with the given start and end location.
+ ///
+ /// \param StartLoc Starting location of the directive kind.
+ /// \param EndLoc Ending location of the directive.
+ /// \param CollapsedNum Number of collapsed nested loops.
+ ///
+ OMPTargetTeamsGenericLoopDirective(SourceLocation StartLoc,
+ SourceLocation EndLoc,
+ unsigned CollapsedNum)
+ : OMPLoopDirective(OMPTargetTeamsGenericLoopDirectiveClass,
+ llvm::omp::OMPD_target_teams_loop, StartLoc, EndLoc,
+ CollapsedNum) {}
+
+ /// Build an empty directive.
+ ///
+ /// \param CollapsedNum Number of collapsed nested loops.
+ ///
+ explicit OMPTargetTeamsGenericLoopDirective(unsigned CollapsedNum)
+ : OMPLoopDirective(OMPTargetTeamsGenericLoopDirectiveClass,
+ llvm::omp::OMPD_target_teams_loop, SourceLocation(),
+ SourceLocation(), CollapsedNum) {}
+
+public:
+ /// Creates directive with a list of \p Clauses.
+ ///
+ /// \param C AST context.
+ /// \param StartLoc Starting location of the directive kind.
+ /// \param EndLoc Ending Location of the directive.
+ /// \param CollapsedNum Number of collapsed loops.
+ /// \param Clauses List of clauses.
+ /// \param AssociatedStmt Statement, associated with the directive.
+ /// \param Exprs Helper expressions for CodeGen.
+ ///
+ static OMPTargetTeamsGenericLoopDirective *
+ Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
+ unsigned CollapsedNum, ArrayRef<OMPClause *> Clauses,
+ Stmt *AssociatedStmt, const HelperExprs &Exprs);
+
+ /// Creates an empty directive with the place
+ /// for \a NumClauses clauses.
+ ///
+ /// \param C AST context.
+ /// \param CollapsedNum Number of collapsed nested loops.
+ /// \param NumClauses Number of clauses.
+ ///
+ static OMPTargetTeamsGenericLoopDirective *CreateEmpty(const ASTContext &C,
+ unsigned NumClauses,
+ unsigned CollapsedNum,
+ EmptyShell);
+
+ static bool classof(const Stmt *T) {
+ return T->getStmtClass() == OMPTargetTeamsGenericLoopDirectiveClass;
+ }
+};
} // end namespace clang
#endif
diff --git a/clang/include/clang/Basic/StmtNodes.td b/clang/include/clang/Basic/StmtNodes.td
index f0de74d5781a2..971970be5182c 100644
--- a/clang/include/clang/Basic/StmtNodes.td
+++ b/clang/include/clang/Basic/StmtNodes.td
@@ -284,3 +284,4 @@ def OMPDispatchDirective : StmtNode<OMPExecutableDirective>;
def OMPMaskedDirective : StmtNode<OMPExecutableDirective>;
def OMPGenericLoopDirective : StmtNode<OMPLoopDirective>;
def OMPTeamsGenericLoopDirective : StmtNode<OMPLoopDirective>;
+def OMPTargetTeamsGenericLoopDirective : StmtNode<OMPLoopDirective>;
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 32a8ffdc8df92..abaf760f18c63 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -10909,6 +10909,11 @@ class Sema final {
StmtResult ActOnOpenMPTeamsGenericLoopDirective(
ArrayRef<OMPClause *> Clauses, Stmt *AStmt, SourceLocation StartLoc,
SourceLocation EndLoc, VarsWithInheritedDSAType &VarsWithImplicitDSA);
+ /// Called on well-formed '\#pragma omp target teams loop' after parsing of
+ /// the associated statement.
+ StmtResult ActOnOpenMPTargetTeamsGenericLoopDirective(
+ ArrayRef<OMPClause *> Clauses, Stmt *AStmt, SourceLocation StartLoc,
+ SourceLocation EndLoc, VarsWithInheritedDSAType &VarsWithImplicitDSA);
/// Called on well-formed '\#pragma omp cancellation point'.
StmtResult
ActOnOpenMPCancellationPointDirective(SourceLocation StartLoc,
diff --git a/clang/include/clang/Serialization/ASTBitCodes.h b/clang/include/clang/Serialization/ASTBitCodes.h
index 3cb3dc622d38c..346514f23c7e3 100644
--- a/clang/include/clang/Serialization/ASTBitCodes.h
+++ b/clang/include/clang/Serialization/ASTBitCodes.h
@@ -1962,6 +1962,7 @@ enum StmtCode {
STMT_OMP_MASKED_DIRECTIVE,
STMT_OMP_GENERIC_LOOP_DIRECTIVE,
STMT_OMP_TEAMS_GENERIC_LOOP_DIRECTIVE,
+ STMT_OMP_TARGET_TEAMS_GENERIC_LOOP_DIRECTIVE,
EXPR_OMP_ARRAY_SECTION,
EXPR_OMP_ARRAY_SHAPING,
EXPR_OMP_ITERATOR,
diff --git a/clang/lib/AST/StmtOpenMP.cpp b/clang/lib/AST/StmtOpenMP.cpp
index 5404dfc6f1592..a7dcc9da3b0b9 100644
--- a/clang/lib/AST/StmtOpenMP.cpp
+++ b/clang/lib/AST/StmtOpenMP.cpp
@@ -2180,3 +2180,48 @@ OMPTeamsGenericLoopDirective::CreateEmpty(const ASTContext &C,
C, NumClauses, /*HasAssociatedStmt=*/true,
numLoopChildren(CollapsedNum, OMPD_teams_loop), CollapsedNum);
}
+
+OMPTargetTeamsGenericLoopDirective *OMPTargetTeamsGenericLoopDirective::Create(
+ const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
+ unsigned CollapsedNum, ArrayRef<OMPClause *> Clauses, Stmt *AssociatedStmt,
+ const HelperExprs &Exprs) {
+ auto *Dir = createDirective<OMPTargetTeamsGenericLoopDirective>(
+ C, Clauses, AssociatedStmt,
+ numLoopChildren(CollapsedNum, OMPD_target_teams_loop), StartLoc, EndLoc,
+ CollapsedNum);
+ Dir->setIterationVariable(Exprs.IterationVarRef);
+ Dir->setLastIteration(Exprs.LastIteration);
+ Dir->setCalcLastIteration(Exprs.CalcLastIteration);
+ Dir->setPreCond(Exprs.PreCond);
+ Dir->setCond(Exprs.Cond);
+ Dir->setInit(Exprs.Init);
+ Dir->setInc(Exprs.Inc);
+ Dir->setIsLastIterVariable(Exprs.IL);
+ Dir->setLowerBoundVariable(Exprs.LB);
+ Dir->setUpperBoundVariable(Exprs.UB);
+ Dir->setStrideVariable(Exprs.ST);
+ Dir->setEnsureUpperBound(Exprs.EUB);
+ Dir->setNextLowerBound(Exprs.NLB);
+ Dir->setNextUpperBound(Exprs.NUB);
+ Dir->setNumIterations(Exprs.NumIterations);
+ Dir->setCounters(Exprs.Counters);
+ Dir->setPrivateCounters(Exprs.PrivateCounters);
+ Dir->setInits(Exprs.Inits);
+ Dir->setUpdates(Exprs.Updates);
+ Dir->setFinals(Exprs.Finals);
+ Dir->setDependentCounters(Exprs.DependentCounters);
+ Dir->setDependentInits(Exprs.DependentInits);
+ Dir->setFinalsConditions(Exprs.FinalsConditions);
+ Dir->setPreInits(Exprs.PreInits);
+ return Dir;
+}
+
+OMPTargetTeamsGenericLoopDirective *
+OMPTargetTeamsGenericLoopDirective::CreateEmpty(const ASTContext &C,
+ unsigned NumClauses,
+ unsigned CollapsedNum,
+ EmptyShell) {
+ return createEmptyDirective<OMPTargetTeamsGenericLoopDirective>(
+ C, NumClauses, /*HasAssociatedStmt=*/true,
+ numLoopChildren(CollapsedNum, OMPD_target_teams_loop), CollapsedNum);
+}
diff --git a/clang/lib/AST/StmtPrinter.cpp b/clang/lib/AST/StmtPrinter.cpp
index bc968cef180ee..4733e3affade4 100644
--- a/clang/lib/AST/StmtPrinter.cpp
+++ b/clang/lib/AST/StmtPrinter.cpp
@@ -1011,6 +1011,12 @@ void StmtPrinter::VisitOMPTeamsGenericLoopDirective(
PrintOMPExecutableDirective(Node);
}
+void StmtPrinter::VisitOMPTargetTeamsGenericLoopDirective(
+ OMPTargetTeamsGenericLoopDirective *Node) {
+ Indent() << "#pragma omp target teams loop";
+ PrintOMPExecutableDirective(Node);
+}
+
//===----------------------------------------------------------------------===//
// Expr printing methods.
//===----------------------------------------------------------------------===//
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index 222970965b4ac..9b765c0bd78e1 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -1212,6 +1212,11 @@ void StmtProfiler::VisitOMPTeamsGenericLoopDirective(
VisitOMPLoopDirective(S);
}
+void StmtProfiler::VisitOMPTargetTeamsGenericLoopDirective(
+ const OMPTargetTeamsGenericLoopDirective *S) {
+ VisitOMPLoopDirective(S);
+}
+
void StmtProfiler::VisitExpr(const Expr *S) {
VisitStmt(S);
}
diff --git a/clang/lib/Basic/OpenMPKinds.cpp b/clang/lib/Basic/OpenMPKinds.cpp
index 52facef6b16d1..505f6aedfe093 100644
--- a/clang/lib/Basic/OpenMPKinds.cpp
+++ b/clang/lib/Basic/OpenMPKinds.cpp
@@ -495,7 +495,8 @@ bool clang::isOpenMPLoopDirective(OpenMPDirectiveKind DKind) {
DKind == OMPD_target_teams_distribute_parallel_for ||
DKind == OMPD_target_teams_distribute_parallel_for_simd ||
DKind == OMPD_target_teams_distribute_simd || DKind == OMPD_tile ||
- DKind == OMPD_unroll || DKind == OMPD_loop || DKind == OMPD_teams_loop;
+ DKind == OMPD_unroll || DKind == OMPD_loop ||
+ DKind == OMPD_teams_loop || DKind == OMPD_target_teams_loop;
}
bool clang::isOpenMPWorksharingDirective(OpenMPDirectiveKind DKind) {
@@ -543,7 +544,8 @@ bool clang::isOpenMPTargetExecutionDirective(OpenMPDirectiveKind DKind) {
DKind == OMPD_target_teams || DKind == OMPD_target_teams_distribute ||
DKind == OMPD_target_teams_distribute_parallel_for ||
DKind == OMPD_target_teams_distribute_parallel_for_simd ||
- DKind == OMPD_target_teams_distribute_simd;
+ DKind == OMPD_target_teams_distribute_simd ||
+ DKind == OMPD_target_teams_loop;
}
bool clang::isOpenMPTargetDataManagementDirective(OpenMPDirectiveKind DKind) {
@@ -560,11 +562,12 @@ bool clang::isOpenMPNestingTeamsDirective(OpenMPDirectiveKind DKind) {
}
bool clang::isOpenMPTeamsDirective(OpenMPDirectiveKind DKind) {
- return isOpenMPNestingTeamsDirective(DKind) ||
- DKind == OMPD_target_teams || DKind == OMPD_target_teams_distribute ||
+ return isOpenMPNestingTeamsDirective(DKind) || DKind == OMPD_target_teams ||
+ DKind == OMPD_target_teams_distribute ||
DKind == OMPD_target_teams_distribute_parallel_for ||
DKind == OMPD_target_teams_distribute_parallel_for_simd ||
- DKind == OMPD_target_teams_distribute_simd;
+ DKind == OMPD_target_teams_distribute_simd ||
+ DKind == OMPD_target_teams_loop;
}
bool clang::isOpenMPSimdDirective(OpenMPDirectiveKind DKind) {
@@ -600,7 +603,8 @@ bool clang::isOpenMPDistributeDirective(OpenMPDirectiveKind Kind) {
}
bool clang::isOpenMPGenericLoopDirective(OpenMPDirectiveKind Kind) {
- return Kind == OMPD_loop || Kind == OMPD_teams_loop;
+ return Kind == OMPD_loop || Kind == OMPD_teams_loop ||
+ Kind == OMPD_target_teams_loop;
}
bool clang::isOpenMPPrivate(OpenMPClauseKind Kind) {
@@ -651,6 +655,7 @@ void clang::getOpenMPCaptureRegions(
case OMPD_target_teams:
case OMPD_target_teams_distribute:
case OMPD_target_teams_distribute_simd:
+ case OMPD_target_teams_loop:
CaptureRegions.push_back(OMPD_task);
CaptureRegions.push_back(OMPD_target);
CaptureRegions.push_back(OMPD_teams);
diff --git a/clang/lib/CodeGen/CGStmt.cpp b/clang/lib/CodeGen/CGStmt.cpp
index ab04e22aa092f..16a881a19bd8b 100644
--- a/clang/lib/CodeGen/CGStmt.cpp
+++ b/clang/lib/CodeGen/CGStmt.cpp
@@ -399,6 +399,9 @@ void CodeGenFunction::EmitStmt(const Stmt *S, ArrayRef<const Attr *> Attrs) {
case Stmt::OMPTeamsGenericLoopDirectiveClass:
llvm_unreachable("teams loop directive not supported yet.");
break;
+ case Stmt::OMPTargetTeamsGenericLoopDirectiveClass:
+ llvm_unreachable("target teams loop directive not supported yet.");
+ break;
}
}
diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp
index 5ab6cd76aed62..02588cf1e05a3 100644
--- a/clang/lib/Parse/ParseOpenMP.cpp
+++ b/clang/lib/Parse/ParseOpenMP.cpp
@@ -166,6 +166,7 @@ static OpenMPDirectiveKindExWrapper parseOpenMPDirectiveKind(Parser &P) {
{OMPD_teams, OMPD_loop, OMPD_teams_loop},
{OMPD_target, OMPD_teams, OMPD_target_teams},
{OMPD_target_teams, OMPD_distribute, OMPD_target_teams_distribute},
+ {OMPD_target_teams, OMPD_loop, OMPD_target_teams_loop},
{OMPD_target_teams_distribute, OMPD_parallel,
OMPD_target_teams_distribute_parallel},
{OMPD_target_teams_distribute, OMPD_simd,
@@ -2401,6 +2402,7 @@ Parser::DeclGroupPtrTy Parser::ParseOpenMPDeclarativeDirectiveWithExtDecl(
case OMPD_metadirective:
case OMPD_loop:
case OMPD_teams_loop:
+ case OMPD_target_teams_loop:
Diag(Tok, diag::err_omp_unexpected_directive)
<< 1 << getOpenMPDirectiveName(DKind);
break;
@@ -2757,6 +2759,7 @@ StmtResult Parser::ParseOpenMPDeclarativeOrExecutableDirective(
case OMPD_target_parallel_for:
case OMPD_loop:
case OMPD_teams_loop:
+ case OMPD_target_teams_loop:
case OMPD_taskloop:
case OMPD_taskloop_simd:
case OMPD_master_taskloop:
diff --git a/clang/lib/Sema/SemaExceptionSpec.cpp b/clang/lib/Sema/SemaExceptionSpec.cpp
index e72495c2413c0..527c4ba7ecf9e 100644
--- a/clang/lib/Sema/SemaExceptionSpec.cpp
+++ b/clang/lib/Sema/SemaExceptionSpec.cpp
@@ -1500,6 +1500,7 @@ CanThrowResult Sema::canThrow(const Stmt *S) {
case Stmt::OMPMetaDirectiveClass:
case Stmt::OMPGenericLoopDirectiveClass:
case Stmt::OMPTeamsGenericLoopDirectiveClass:
+ case Stmt::OMPTargetTeamsGenericLoopDirectiveClass:
case Stmt::ReturnStmtClass:
case Stmt::SEHExceptStmtClass:
case Stmt::SEHFinallyStmtClass:
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 8e8535a0f98f6..cfd8b60d9f03c 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -3954,6 +3954,7 @@ void Sema::ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, Scope *CurScope) {
case OMPD_target_parallel:
case OMPD_target_parallel_for:
case OMPD_target_parallel_for_simd:
+ case OMPD_target_teams_loop:
case OMPD_target_teams_distribute:
case OMPD_target_teams_distribute_simd: {
QualType KmpInt32Ty = Context.getIntTypeForBitwidth(32, 1).withConst();
@@ -6315,6 +6316,10 @@ StmtResult Sema::ActOnOpenMPExecutableDirective(
Res = ActOnOpenMPTeamsGenericLoopDirective(
ClausesWithImplicit, AStmt, StartLoc, EndLoc, VarsWithInheritedDSA);
break;
+ case OMPD_target_teams_loop:
+ Res = ActOnOpenMPTargetTeamsGenericLoopDirective(
+ ClausesWithImplicit, AStmt, StartLoc, EndLoc, VarsWithInheritedDSA);
+ break;
case OMPD_declare_target:
case OMPD_end_declare_target:
case OMPD_threadprivate:
@@ -10258,6 +10263,55 @@ StmtResult Sema::ActOnOpenMPTeamsGenericLoopDirective(
Context, StartLoc, EndLoc, NestedLoopCount, Clauses, AStmt, B);
}
+StmtResult Sema::ActOnOpenMPTargetTeamsGenericLoopDirective(
+ ArrayRef<OMPClause *> Clauses, Stmt *AStmt, SourceLocation StartLoc,
+ SourceLocation EndLoc, VarsWithInheritedDSAType &VarsWithImplicitDSA) {
+ if (!AStmt)
+ return StmtError();
+
+ // OpenMP 5.1 [2.11.7, loop construct, Restrictions]
+ // A list item may not appear in a lastprivate clause unless it is the
+ // loop iteration variable of a loop that is associated with the construct.
+ if (checkGenericLoopLastprivate(*this, Clauses, OMPD_target_teams_loop,
+ DSAStack))
+ return StmtError();
+
+ auto *CS = cast<CapturedStmt>(AStmt);
+ // 1.2.2 OpenMP Language Terminology
+ // Structured block - An executable statement with a single entry at the
+ // top and a single exit at the bottom.
+ // The point of exit cannot be a branch out of the structured block.
+ // longjmp() and throw() must not violate the entry/exit criteria.
+ CS->getCapturedDecl()->setNothrow();
+ for (int ThisCaptureLevel = getOpenMPCaptureLevels(OMPD_target_teams_loop);
+ ThisCaptureLevel > 1; --ThisCaptureLevel) {
+ CS = cast<CapturedStmt>(CS->getCapturedStmt());
+ // 1.2.2 OpenMP Language Terminology
+ // Structured block - An executable statement with a single entry at the
+ // top and a single exit at the bottom.
+ // The point of exit cannot be a branch out of the structured block.
+ // longjmp() and throw() must not violate the entry/exit criteria.
+ CS->getCapturedDecl()->setNothrow();
+ }
+
+ OMPLoopDirective::HelperExprs B;
+ // In presence of clause 'collapse', it will define the nested loops number.
+ unsigned NestedLoopCount =
+ checkOpenMPLoop(OMPD_target_teams_loop, getCollapseNumberExpr(Clauses),
+ /*OrderedLoopCountExpr=*/nullptr, CS, *this, *DSAStack,
+ VarsWithImplicitDSA, B);
+ if (NestedLoopCount == 0)
+ return StmtError();
+
+ assert((CurContext->isDependentContext() || B.builtAll()) &&
+ "omp loop exprs were not built");
+
+ setFunctionHasBranchProtectedScope();
+
+ return OMPTargetTeamsGenericLoopDirective::Create(
+ Context, StartLoc, EndLoc, NestedLoopCount, Clauses, AStmt, B);
+}
+
StmtResult Sema::ActOnOpenMPSingleDirective(ArrayRef<OMPClause *> Clauses,
Stmt *AStmt,
SourceLocation StartLoc,
@@ -14568,6 +14622,7 @@ static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
case OMPD_target:
case OMPD_target_teams:
case OMPD_target_teams_distribute:
+ case OMPD_target_teams_loop:
case OMPD_distribute_parallel_for:
case OMPD_task:
case OMPD_taskloop:
@@ -14678,6 +14733,7 @@ static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
case OMPD_end_declare_target:
case OMPD_loop:
case OMPD_teams_loop:
+ case OMPD_target_teams_loop:
case OMPD_teams:
case OMPD_simd:
case OMPD_tile:
@@ -14712,6 +14768,7 @@ static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
case OMPD_target_teams_distribute_simd:
case OMPD_target_teams_distribute_parallel_for:
case OMPD_target_teams_distribute_parallel_for_simd:
+ case OMPD_target_teams_loop:
CaptureRegion = OMPD_target;
break;
case OMPD_teams_distribute_parallel_for:
@@ -14795,6 +14852,7 @@ static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
case OMPD_target_teams_distribute_simd:
case OMPD_target_teams_distribute_parallel_for:
case OMPD_target_teams_distribute_parallel_for_simd:
+ case OMPD_target_teams_loop:
CaptureRegion = OMPD_target;
break;
case OMPD_teams_distribute_parallel_for:
@@ -14931,6 +14989,7 @@ static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
case OMPD_end_declare_target:
case OMPD_loop:
case OMPD_teams_loop:
+ case OMPD_target_teams_loop:
case OMPD_simd:
case OMPD_tile:
case OMPD_unroll:
@@ -15014,6 +15073,7 @@ static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
case OMPD_end_declare_target:
case OMPD_loop:
case OMPD_teams_loop:
+ case OMPD_target_teams_loop:
case OMPD_simd:
case OMPD_tile:
case OMPD_unroll:
@@ -15052,6 +15112,7 @@ static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
case OMPD_target_parallel_for_simd:
case OMPD_target_teams_distribute_parallel_for:
case OMPD_target_teams_distribute_parallel_for_simd:
+ case OMPD_target_teams_loop:
case OMPD_dispatch:
CaptureRegion = OMPD_task;
break;
@@ -15183,6 +15244,7 @@ static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
case OMPD_end_declare_target:
case OMPD_loop:
case OMPD_teams_loop:
+ case OMPD_target_teams_loop:
case OMPD_simd:
case OMPD_tile:
case OMPD_unroll:
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index 01ab106e2fee5..31d7c5a13fc1f 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -9260,6 +9260,17 @@ StmtResult TreeTransform<Derived>::TransformOMPTeamsGenericLoopDirective(
return Res;
}
+template <typename Derived>
+StmtResult TreeTransform<Derived>::TransformOMPTargetTeamsGenericLoopDirective(
+ OMPTargetTeamsGenericLoopDirective *D) {
+ DeclarationNameInfo DirName;
+ getDerived().getSema().StartOpenMPDSABlock(OMPD_target_teams_loop, DirName,
+ nullptr, D->getBeginLoc());
+ StmtResult Res = getDerived().TransformOMPExecutableDirective(D);
+ getDerived().getSema().EndOpenMPDSABlock(Res.get());
+ return Res;
+}
+
//===----------------------------------------------------------------------===//
// OpenMP clause transformation
//===----------------------------------------------------------------------===//
diff --git a/clang/lib/Serialization/ASTReaderStmt.cpp b/clang/lib/Serialization/ASTReaderStmt.cpp
index e6589fe0701ba..028f8a5d4771f 100644
--- a/clang/lib/Serialization/ASTReaderStmt.cpp
+++ b/clang/lib/Serialization/ASTReaderStmt.cpp
@@ -2643,6 +2643,11 @@ void ASTStmtReader::VisitOMPTeamsGenericLoopDirective(
VisitOMPLoopDirective(D);
}
+void ASTStmtReader::VisitOMPTargetTeamsGenericLoopDirective(
+ OMPTargetTeamsGenericLoopDirective *D) {
+ VisitOMPLoopDirective(D);
+}
+
//===----------------------------------------------------------------------===//
// ASTReader Implementation
//===----------------------------------------------------------------------===//
@@ -3606,6 +3611,14 @@ Stmt *ASTReader::ReadStmtFromStream(ModuleFile &F) {
break;
}
+ case STMT_OMP_TARGET_TEAMS_GENERIC_LOOP_DIRECTIVE: {
+ unsigned CollapsedNum = Record[ASTStmtReader::NumStmtFields];
+ unsigned NumClauses = Record[ASTStmtReader::NumStmtFields + 1];
+ S = OMPTargetTeamsGenericLoopDirective::CreateEmpty(Context, NumClauses,
+ CollapsedNum, Empty);
+ break;
+ }
+
case EXPR_CXX_OPERATOR_CALL:
S = CXXOperatorCallExpr::CreateEmpty(
Context, /*NumArgs=*/Record[ASTStmtReader::NumExprFields],
diff --git a/clang/lib/Serialization/ASTWriterStmt.cpp b/clang/lib/Serialization/ASTWriterStmt.cpp
index f675ac78d62f2..8a616bffb2583 100644
--- a/clang/lib/Serialization/ASTWriterStmt.cpp
+++ b/clang/lib/Serialization/ASTWriterStmt.cpp
@@ -2601,6 +2601,12 @@ void ASTStmtWriter::VisitOMPTeamsGenericLoopDirective(
Code = serialization::STMT_OMP_TEAMS_GENERIC_LOOP_DIRECTIVE;
}
+void ASTStmtWriter::VisitOMPTargetTeamsGenericLoopDirective(
+ OMPTargetTeamsGenericLoopDirective *D) {
+ VisitOMPLoopDirective(D);
+ Code = serialization::STMT_OMP_TARGET_TEAMS_GENERIC_LOOP_DIRECTIVE;
+}
+
//===----------------------------------------------------------------------===//
// ASTWriter Implementation
//===----------------------------------------------------------------------===//
diff --git a/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp b/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
index 00275f288a0a4..0f1f8cef5d5ca 100644
--- a/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
+++ b/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
@@ -1298,6 +1298,7 @@ void ExprEngine::Visit(const Stmt *S, ExplodedNode *Pred,
case Stmt::OMPMaskedDirectiveClass:
case Stmt::OMPGenericLoopDirectiveClass:
case Stmt::OMPTeamsGenericLoopDirectiveClass:
+ case Stmt::OMPTargetTeamsGenericLoopDirectiveClass:
case Stmt::CapturedStmtClass:
case Stmt::OMPUnrollDirectiveClass:
case Stmt::OMPMetaDirectiveClass: {
diff --git a/clang/test/Analysis/cfg-openmp.cpp b/clang/test/Analysis/cfg-openmp.cpp
index 291dac7fa5d1b..4939b49dd69bf 100644
--- a/clang/test/Analysis/cfg-openmp.cpp
+++ b/clang/test/Analysis/cfg-openmp.cpp
@@ -701,3 +701,27 @@ void teamsloop(int argc) {
for (int i = 0; i < 10; ++i)
argc = x;
}
+
+// CHECK-LABEL: void targetteamsloop(int argc)
+void targetteamsloop(int argc) {
+ int x, cond, fp, rd, map;
+// CHECK-DAG: [B3]
+// CHECK-DAG: [[#TTDB:]]: x
+// CHECK-DAG: [[#TTDB+1]]: [B3.[[#TTDB]]] (ImplicitCastExpr, LValueToRValue, int)
+// CHECK-DAG: [[#TTDB+2]]: argc
+// CHECK-DAG: [[#TTDB+3]]: [B3.[[#TTDB+2]]] = [B3.[[#TTDB+1]]]
+// CHECK-DAG: [B1]
+// CHECK-DAG: [[#TTD:]]: cond
+// CHECK-DAG: [[#TTD+1]]: [B1.[[#TTD]]] (ImplicitCastExpr, LValueToRValue, int)
+// CHECK-DAG: [[#TTD+2]]: [B1.[[#TTD+1]]] (ImplicitCastExpr, IntegralToBoolean, _Bool)
+// CHECK-DAG: [[#TTD+3]]: fp
+// CHECK-DAG: [[#TTD+4]]: rd
+// CHECK-DAG: [[#TTD+5]]: [B3.[[#TTDB+2]]]
+// CHECK-DAG: [[#TTD+6]]: [B3.[[#TTDB]]]
+// CHECK-DAG: [[#TTD+7]]: #pragma omp target teams loop if(cond) firstprivate(fp) reduction(+: rd) map(alloc: map)
+// CHECK-DAG: for (int i = 0;
+// CHECK-DAG: [B3.[[#TTDB+3]]];
+#pragma omp target teams loop if(cond) firstprivate(fp) reduction(+:rd) map(alloc:map)
+ for (int i = 0; i < 10; ++i)
+ argc = x;
+}
diff --git a/clang/test/OpenMP/target_teams_generic_loop_ast_print.cpp b/clang/test/OpenMP/target_teams_generic_loop_ast_print.cpp
new file mode 100644
index 0000000000000..92094faa3fe08
--- /dev/null
+++ b/clang/test/OpenMP/target_teams_generic_loop_ast_print.cpp
@@ -0,0 +1,123 @@
+// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -fopenmp -fopenmp-version=51 \
+// RUN: -fsyntax-only -verify %s
+
+// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -fopenmp -fopenmp-version=51 \
+// RUN: -ast-print %s | FileCheck %s
+
+// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -fopenmp -fopenmp-version=51 \
+// RUN: -emit-pch -o %t %s
+
+// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -fopenmp -fopenmp-version=51 \
+// RUN: -include-pch %t -ast-print %s | FileCheck %s
+
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+typedef void **omp_allocator_handle_t;
+extern const omp_allocator_handle_t omp_null_allocator;
+extern const omp_allocator_handle_t omp_default_mem_alloc;
+extern const omp_allocator_handle_t omp_large_cap_mem_alloc;
+extern const omp_allocator_handle_t omp_const_mem_alloc;
+extern const omp_allocator_handle_t omp_high_bw_mem_alloc;
+extern const omp_allocator_handle_t omp_low_lat_mem_alloc;
+extern const omp_allocator_handle_t omp_cgroup_mem_alloc;
+extern const omp_allocator_handle_t omp_pteam_mem_alloc;
+extern const omp_allocator_handle_t omp_thread_mem_alloc;
+
+//CHECK: template <typename T, int C, int D> void templ_foo(T t) {
+//CHECK: T j, z;
+//CHECK: #pragma omp target teams loop device(D) collapse(C) reduction(+: z) lastprivate(j) bind(thread) num_teams(C + 2)
+//CHECK: for (T i = 0; i < t; ++i)
+//CHECK: for (j = 0; j < t; ++j)
+//CHECK: z += i + j;
+//CHECK: }
+
+//CHECK: template<> void templ_foo<int, 2, 0>(int t) {
+//CHECK: int j, z;
+//CHECK: #pragma omp target teams loop device(0) collapse(2) reduction(+: z) lastprivate(j) bind(thread) num_teams(2 + 2)
+//CHECK: for (int i = 0; i < t; ++i)
+//CHECK: for (j = 0; j < t; ++j)
+//CHECK: z += i + j;
+//CHECK: }
+template <typename T, int C, int D>
+void templ_foo(T t) {
+
+ T j,z;
+ #pragma omp target teams loop device(D) collapse(C) reduction(+:z) lastprivate(j) bind(thread) num_teams(C+2)
+ for (T i = 0; i<t; ++i)
+ for (j = 0; j<t; ++j)
+ z += i+j;
+}
+
+
+//CHECK: void test() {
+void test() {
+ constexpr int N = 100;
+ float MTX[N][N];
+ int aaa[1000];
+
+ //CHECK: #pragma omp target teams loop map(tofrom: MTX)
+ #pragma omp target teams loop map(MTX)
+ for (auto j = 0; j < N; ++j) {
+ MTX[0][j] = 0;
+ }
+
+ int j, z, z1;
+ //CHECK: #pragma omp target teams loop collapse(2) private(z) lastprivate(j) order(concurrent) reduction(+: z1) bind(parallel)
+ #pragma omp target teams loop collapse(2) private(z) lastprivate(j) \
+ order(concurrent) reduction(+:z1) bind(parallel)
+ for (auto i = 0; i < N; ++i) {
+ for (j = 0; j < N; ++j) {
+ z = i+j;
+ MTX[i][j] = z;
+ z1 += z;
+ }
+ }
+
+ //CHECK: #pragma omp target teams loop bind(teams) num_teams(16) thread_limit(8) default(none)
+ #pragma omp target teams loop bind(teams) num_teams(16) thread_limit(8) default(none)
+ for (auto i = 0; i < N; ++i) { }
+
+ int pr;
+ int zzz;
+ //CHECK: #pragma omp target teams loop private(zzz) uses_allocators(omp_default_mem_alloc) allocate(omp_default_mem_alloc: zzz) if(1) device(0) map(tofrom: pr)
+ #pragma omp target teams loop private(zzz) uses_allocators(omp_default_mem_alloc) allocate(omp_default_mem_alloc:zzz) if(1) device(0) map(tofrom:pr)
+ for (int i=0; i<1000; ++i) {
+ zzz = i + 1;
+ pr = 33;
+ }
+
+ int fpr = 10;
+ int k;
+ int s = 20;
+ //CHECK: #pragma omp target teams loop bind(teams) private(pr) firstprivate(fpr) shared(s) allocate(k) reduction(+: k)
+ #pragma omp target teams loop bind(teams) private(pr) firstprivate(fpr) \
+ shared(s) allocate(k) reduction(+:k)
+ for (auto i = 0; i < N; ++i) {
+ pr = i + fpr + s;
+ }
+
+ short y = 3;
+ //CHECK: #pragma omp target teams loop map(tofrom: y) depend(out : y)
+ #pragma omp target teams loop map(tofrom:y) depend(out:y)
+ for (int i=0; i<10; ++i) {
+ y = 3+i;
+ }
+}
+
+//CHECK: void nobindingfunc() {
+void nobindingfunc()
+{
+ //CHECK: #pragma omp target teams loop
+ #pragma omp target teams loop
+ for (int i=0; i<10; ++i) { }
+}
+
+void bar()
+{
+ templ_foo<int,2,0>(8);
+}
+
+#endif // HEADER
diff --git a/clang/test/OpenMP/target_teams_generic_loop_messages.cpp b/clang/test/OpenMP/target_teams_generic_loop_messages.cpp
new file mode 100644
index 0000000000000..e272e93d3c190
--- /dev/null
+++ b/clang/test/OpenMP/target_teams_generic_loop_messages.cpp
@@ -0,0 +1,152 @@
+// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -verify -fopenmp \
+// RUN: -fopenmp-version=51 -Wuninitialized %s
+
+void foo()
+{
+ int i,j,k;
+ int z;
+
+ // expected-error at +2 {{statement after '#pragma omp target teams loop' must be a for loop}}
+ #pragma omp target teams loop bind(thread)
+ i = 0;
+
+ // OpenMP 5.1 [2.22 Nesting of regions]
+ //
+ // A barrier region may not be closely nested inside a worksharing, loop,
+ // task, taskloop, critical, ordered, atomic, or masked region.
+
+ // expected-error at +3 {{region cannot be closely nested inside 'target teams loop' region}}
+ #pragma omp target teams loop bind(thread)
+ for (i=0; i<1000; ++i) {
+ #pragma omp barrier
+ }
+
+ // A masked region may not be closely nested inside a worksharing, loop,
+ // atomic, task, or taskloop region.
+
+ // expected-error at +3 {{region cannot be closely nested inside 'target teams loop' region}}
+ #pragma omp target teams loop bind(thread)
+ for (i=0; i<1000; ++i) {
+ #pragma omp masked filter(2)
+ { }
+ }
+
+ // An ordered region that corresponds to an ordered construct without any
+ // clause or with the threads or depend clause may not be closely nested
+ // inside a critical, ordered, loop, atomic, task, or taskloop region.
+
+ // expected-error at +3 {{region cannot be closely nested inside 'target teams loop' region; perhaps you forget to enclose 'omp ordered' directive into a for or a parallel for region with 'ordered' clause?}}
+ #pragma omp target teams loop bind(thread)
+ for (i=0; i<1000; ++i) {
+ #pragma omp ordered
+ { }
+ }
+
+ // expected-error at +3 {{region cannot be closely nested inside 'target teams loop' region; perhaps you forget to enclose 'omp ordered' directive into a for or a parallel for region with 'ordered' clause?}}
+ #pragma omp target teams loop bind(thread)
+ for (i=0; i<1000; ++i) {
+ #pragma omp ordered threads
+ { }
+ }
+
+ // expected-error at +3 {{region cannot be closely nested inside 'target teams loop' region; perhaps you forget to enclose 'omp ordered' directive into a for or a parallel for region with 'ordered' clause?}}
+ #pragma omp target teams loop bind(thread)
+ for (i=0; i<1000; ++i) {
+ #pragma omp ordered depend(source)
+ }
+
+ // bind clause
+
+ // expected-error at +1 {{directive '#pragma omp target teams loop' cannot contain more than one 'bind' clause}}
+ #pragma omp target teams loop bind(thread) bind(thread)
+ for (i=0; i<1000; ++i) {
+ }
+
+ // expected-error at +1 {{expected 'teams', 'parallel' or 'thread' in OpenMP clause 'bind'}}
+ #pragma omp target teams loop bind(other)
+ for (i=0; i<1000; ++i) {
+ }
+
+ // collapse clause
+
+ // expected-error at +4 {{expected 2 for loops after '#pragma omp target teams loop', but found only 1}}
+ // expected-note at +1 {{as specified in 'collapse' clause}}
+ #pragma omp target teams loop collapse(2) bind(thread)
+ for (i=0; i<1000; ++i)
+ z = i+11;
+
+ // expected-error at +1 {{directive '#pragma omp target teams loop' cannot contain more than one 'collapse' clause}}
+ #pragma omp target teams loop collapse(2) collapse(2) bind(thread)
+ for (i=0; i<1000; ++i)
+ for (j=0; j<1000; ++j)
+ z = i+j+11;
+
+ // order clause
+
+ // expected-error at +1 {{expected 'concurrent' in OpenMP clause 'order'}}
+ #pragma omp target teams loop order(foo) bind(thread)
+ for (i=0; i<1000; ++i)
+ z = i+11;
+
+ // private clause
+
+ // expected-error at +1 {{use of undeclared identifier 'undef_var'}}
+ #pragma omp target teams loop private(undef_var) bind(thread)
+ for (i=0; i<1000; ++i)
+ z = i+11;
+
+ // lastprivate
+
+ // A list item may not appear in a lastprivate clause unless it is the loop
+ // iteration variable of a loop that is associated with the construct.
+
+ // expected-error at +1 {{only loop iteration variables are allowed in 'lastprivate' clause in 'omp target teams loop' directives}}
+ #pragma omp target teams loop lastprivate(z) bind(thread)
+ for (i=0; i<1000; ++i) {
+ z = i+11;
+ }
+
+ // expected-error at +1 {{only loop iteration variables are allowed in 'lastprivate' clause in 'omp target teams loop' directives}}
+ #pragma omp target teams loop lastprivate(k) collapse(2) bind(thread)
+ for (i=0; i<1000; ++i)
+ for (j=0; j<1000; ++j)
+ for (k=0; k<1000; ++k)
+ z = i+j+k+11;
+
+ // reduction
+
+ // expected-error at +1 {{use of undeclared identifier 'undef_var'}}
+ #pragma omp target teams loop reduction(+:undef_var) bind(thread)
+ for (i=0; i<1000; ++i)
+ z = i+11;
+
+ // nowait
+
+ // expected-error at +1 {{directive '#pragma omp target teams loop' cannot contain more than one 'nowait' clause}}
+ #pragma omp target teams loop nowait nowait
+ for (i=0; i<1000; ++i)
+ z = i+11;
+
+}
+
+template <typename T, int C>
+void templ_test(T t) {
+ T i,z;
+
+ // expected-error at +4 {{expected 2 for loops after '#pragma omp target teams loop', but found only 1}}
+ // expected-note at +1 {{as specified in 'collapse' clause}}
+ #pragma omp target teams loop collapse(C) bind(thread)
+ for (i=0; i<1000; ++i)
+ z = i+11;
+
+ // expected-error at +1 {{only loop iteration variables are allowed in 'lastprivate' clause in 'omp target teams loop' directives}}
+ #pragma omp target teams loop lastprivate(z) bind(thread)
+ for (i=0; i<1000; ++i) {
+ z = i+11;
+ }
+}
+
+void bar()
+{
+ templ_test<int, 2>(16); // expected-note {{in instantiation of function template specialization 'templ_test<int, 2>' requested here}}
+}
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index a01dd418e5d23..b2e4557c4a1ed 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -5728,6 +5728,8 @@ CXString clang_getCursorKindSpelling(enum CXCursorKind Kind) {
return cxstring::createRef("OMPGenericLoopDirective");
case CXCursor_OMPTeamsGenericLoopDirective:
return cxstring::createRef("OMPTeamsGenericLoopDirective");
+ case CXCursor_OMPTargetTeamsGenericLoopDirective:
+ return cxstring::createRef("OMPTargetTeamsGenericLoopDirective");
case CXCursor_OverloadCandidate:
return cxstring::createRef("OverloadCandidate");
case CXCursor_TypeAliasTemplateDecl:
diff --git a/clang/tools/libclang/CXCursor.cpp b/clang/tools/libclang/CXCursor.cpp
index 770a7f811dd82..1bc79765e64dd 100644
--- a/clang/tools/libclang/CXCursor.cpp
+++ b/clang/tools/libclang/CXCursor.cpp
@@ -829,6 +829,9 @@ CXCursor cxcursor::MakeCXCursor(const Stmt *S, const Decl *Parent,
case Stmt::OMPTeamsGenericLoopDirectiveClass:
K = CXCursor_OMPTeamsGenericLoopDirective;
break;
+ case Stmt::OMPTargetTeamsGenericLoopDirectiveClass:
+ K = CXCursor_OMPTargetTeamsGenericLoopDirective;
+ break;
case Stmt::BuiltinBitCastExprClass:
K = CXCursor_BuiltinBitCastExpr;
}
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMP.td b/llvm/include/llvm/Frontend/OpenMP/OMP.td
index 18eb54e5680c5..cbe97c866f5b6 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMP.td
+++ b/llvm/include/llvm/Frontend/OpenMP/OMP.td
@@ -1782,6 +1782,32 @@ def OMP_teams_loop : Directive<"teams loop"> {
VersionedClause<OMPC_ThreadLimit>,
];
}
+def OMP_target_teams_loop : Directive<"target teams loop"> {
+ let allowedClauses = [
+ VersionedClause<OMPC_Allocate>,
+ VersionedClause<OMPC_Depend>,
+ VersionedClause<OMPC_DefaultMap>,
+ VersionedClause<OMPC_Device>,
+ VersionedClause<OMPC_FirstPrivate>,
+ VersionedClause<OMPC_IsDevicePtr>,
+ VersionedClause<OMPC_LastPrivate>,
+ VersionedClause<OMPC_Map>,
+ VersionedClause<OMPC_Private>,
+ VersionedClause<OMPC_Reduction>,
+ VersionedClause<OMPC_Shared>,
+ VersionedClause<OMPC_UsesAllocators, 50>
+ ];
+ let allowedOnceClauses = [
+ VersionedClause<OMPC_Bind, 50>,
+ VersionedClause<OMPC_Collapse>,
+ VersionedClause<OMPC_Default>,
+ VersionedClause<OMPC_If>,
+ VersionedClause<OMPC_NoWait>,
+ VersionedClause<OMPC_NumTeams>,
+ VersionedClause<OMPC_Order>,
+ VersionedClause<OMPC_ThreadLimit>,
+ ];
+}
def OMP_Metadirective : Directive<"metadirective"> {
let allowedClauses = [VersionedClause<OMPC_When>];
let allowedOnceClauses = [VersionedClause<OMPC_Default>];
More information about the llvm-commits
mailing list