[OPENMP] Initial codegen for `target teams distribute parallel for
simd`.
Added host codegen + codegen for devices with default codegen for
`#pragma omp target teams distribute parallel for simd` directive.
llvm-svn: 322515
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 4f45d6d..b49f6a4 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -2266,8 +2266,7 @@
break;
}
case OMPD_distribute_parallel_for_simd:
- case OMPD_distribute_parallel_for:
- case OMPD_target_teams_distribute_parallel_for_simd: {
+ case OMPD_distribute_parallel_for: {
QualType KmpInt32Ty = Context.getIntTypeForBitwidth(32, 1);
QualType KmpInt32PtrTy =
Context.getPointerType(KmpInt32Ty).withConst().withRestrict();
@@ -2282,7 +2281,8 @@
Params);
break;
}
- case OMPD_target_teams_distribute_parallel_for: {
+ case OMPD_target_teams_distribute_parallel_for:
+ case OMPD_target_teams_distribute_parallel_for_simd: {
QualType KmpInt32Ty = Context.getIntTypeForBitwidth(32, 1);
QualType KmpInt32PtrTy =
Context.getPointerType(KmpInt32Ty).withConst().withRestrict();
@@ -7490,7 +7490,6 @@
// 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_distribute_parallel_for);
ThisCaptureLevel > 1; --ThisCaptureLevel) {
@@ -7516,6 +7515,17 @@
assert((CurContext->isDependentContext() || B.builtAll()) &&
"omp target teams distribute parallel for loop exprs were not built");
+ if (!CurContext->isDependentContext()) {
+ // Finalize the clauses that need pre-built expressions for CodeGen.
+ for (auto C : Clauses) {
+ if (auto *LC = dyn_cast<OMPLinearClause>(C))
+ if (FinishOpenMPLinearClause(*LC, cast<DeclRefExpr>(B.IterationVarRef),
+ B.NumIterations, *this, CurScope,
+ DSAStack))
+ return StmtError();
+ }
+ }
+
getCurFunction()->setHasBranchProtectedScope();
return OMPTargetTeamsDistributeParallelForDirective::Create(
Context, StartLoc, EndLoc, NestedLoopCount, Clauses, AStmt, B,
@@ -7536,15 +7546,26 @@
// 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_distribute_parallel_for_simd);
+ 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' with number of loops, it will
// define the nested loops number.
- auto NestedLoopCount = CheckOpenMPLoop(
- OMPD_target_teams_distribute_parallel_for_simd,
- getCollapseNumberExpr(Clauses),
- nullptr /*ordered not a clause on distribute*/, AStmt, *this, *DSAStack,
- VarsWithImplicitDSA, B);
+ auto NestedLoopCount =
+ CheckOpenMPLoop(OMPD_target_teams_distribute_parallel_for_simd,
+ getCollapseNumberExpr(Clauses),
+ nullptr /*ordered not a clause on distribute*/, CS, *this,
+ *DSAStack, VarsWithImplicitDSA, B);
if (NestedLoopCount == 0)
return StmtError();
@@ -9209,7 +9230,7 @@
// A list item cannot appear in both a map clause and a data-sharing
// attribute clause on the same construct
if (CurrDir == OMPD_target || CurrDir == OMPD_target_parallel ||
- CurrDir == OMPD_target_teams ||
+ CurrDir == OMPD_target_teams ||
CurrDir == OMPD_target_teams_distribute ||
CurrDir == OMPD_target_teams_distribute_parallel_for ||
CurrDir == OMPD_target_teams_distribute_parallel_for_simd ||