[llvm-branch-commits] [clang] [clang][OpenMP] Extend no-loop promotion to the split teams loop (PR #224650)
Nicole Aschenbrenner via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Wed Oct 7 05:47:41 PDT 2026
https://github.com/nicebert updated https://github.com/llvm/llvm-project/pull/224650
>From f571d5c3089bef65d6ecc4f300bd1112a4c3e3d0 Mon Sep 17 00:00:00 2001
From: Nicole Aschenbrenner <nicole.aschenbrenner at amd.com>
Date: Fri, 18 Sep 2026 08:20:38 -0500
Subject: [PATCH] [clang][OpenMP] Extend no-loop promotion to the split teams
loop
Widen the no-loop promotion to apply not only to the fused target teams
distribute parallel for but include the same region written as separate
directives, other directives promotable to parallel for, and regions
carrying a schedule clause that asks for a mapping a no-loop kernel
already provides.
Use the SPMD_NO_LOOP kernel tag to carry the promotion decision, based
on whether the kernel can get fully promoted to a no-loop kernel. This
eliminates the previous discrepancy where a clause got the tag but not
the promotion.
---
clang/include/clang/AST/StmtOpenMP.h | 11 +-
clang/lib/AST/StmtOpenMP.cpp | 3 +-
clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp | 66 ++++++++++--
clang/lib/CodeGen/CGStmtOpenMP.cpp | 67 +++++++-----
clang/lib/Sema/SemaOpenMP.cpp | 3 +-
clang/lib/Serialization/ASTReaderStmt.cpp | 1 +
clang/lib/Serialization/ASTWriterStmt.cpp | 1 +
clang/test/OpenMP/target_no_loop.cpp | 119 +++++++++++++++++-----
8 files changed, 206 insertions(+), 65 deletions(-)
diff --git a/clang/include/clang/AST/StmtOpenMP.h b/clang/include/clang/AST/StmtOpenMP.h
index 773e1153ccdfc..f2f89cac0ebda 100644
--- a/clang/include/clang/AST/StmtOpenMP.h
+++ b/clang/include/clang/AST/StmtOpenMP.h
@@ -6641,6 +6641,8 @@ class OMPGenericLoopDirective final : public OMPLoopDirective {
class OMPTeamsGenericLoopDirective final : public OMPLoopDirective {
friend class ASTStmtReader;
friend class OMPExecutableDirective;
+ /// true if loop directive's associated loop can be a parallel for.
+ bool CanBeParallelFor = false;
/// Build directive with the given start and end location.
///
/// \param StartLoc Starting location of the directive kind.
@@ -6662,6 +6664,9 @@ class OMPTeamsGenericLoopDirective final : public OMPLoopDirective {
llvm::omp::OMPD_teams_loop, SourceLocation(),
SourceLocation(), CollapsedNum) {}
+ /// Set whether associated loop can be a parallel for.
+ void setCanBeParallelFor(bool ParFor) { CanBeParallelFor = ParFor; }
+
public:
/// Creates directive with a list of \p Clauses.
///
@@ -6676,7 +6681,7 @@ class OMPTeamsGenericLoopDirective final : public OMPLoopDirective {
static OMPTeamsGenericLoopDirective *
Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
unsigned CollapsedNum, ArrayRef<OMPClause *> Clauses,
- Stmt *AssociatedStmt, const HelperExprs &Exprs);
+ Stmt *AssociatedStmt, const HelperExprs &Exprs, bool CanBeParallelFor);
/// Creates an empty directive with the place
/// for \a NumClauses clauses.
@@ -6690,6 +6695,10 @@ class OMPTeamsGenericLoopDirective final : public OMPLoopDirective {
unsigned CollapsedNum,
EmptyShell);
+ /// Return true if current loop directive's associated loop can be a
+ /// parallel for.
+ bool canBeParallelFor() const { return CanBeParallelFor; }
+
static bool classof(const Stmt *T) {
return T->getStmtClass() == OMPTeamsGenericLoopDirectiveClass;
}
diff --git a/clang/lib/AST/StmtOpenMP.cpp b/clang/lib/AST/StmtOpenMP.cpp
index 555e3454816ce..6a82d0e6f9faa 100644
--- a/clang/lib/AST/StmtOpenMP.cpp
+++ b/clang/lib/AST/StmtOpenMP.cpp
@@ -2619,7 +2619,7 @@ OMPGenericLoopDirective::CreateEmpty(const ASTContext &C, unsigned NumClauses,
OMPTeamsGenericLoopDirective *OMPTeamsGenericLoopDirective::Create(
const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
unsigned CollapsedNum, ArrayRef<OMPClause *> Clauses, Stmt *AssociatedStmt,
- const HelperExprs &Exprs) {
+ const HelperExprs &Exprs, bool CanBeParallelFor) {
auto *Dir = createDirective<OMPTeamsGenericLoopDirective>(
C, Clauses, AssociatedStmt,
numLoopChildren(CollapsedNum, OMPD_teams_loop), StartLoc, EndLoc,
@@ -2661,6 +2661,7 @@ OMPTeamsGenericLoopDirective *OMPTeamsGenericLoopDirective::Create(
Dir->setCombinedNextUpperBound(Exprs.DistCombinedFields.NUB);
Dir->setCombinedDistCond(Exprs.DistCombinedFields.DistCond);
Dir->setCombinedParForInDistCond(Exprs.DistCombinedFields.ParForInDistCond);
+ Dir->setCanBeParallelFor(CanBeParallelFor);
return Dir;
}
diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
index a860aace68b39..383d5fe7e0d83 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
@@ -706,27 +706,73 @@ static bool supportsSPMDExecutionMode(ASTContext &Ctx,
"Unknown programming model for OpenMP directive on NVPTX target.");
}
+static bool isNoLoopCompatibleSchedule(ASTContext &Ctx,
+ const OMPLoopDirective &LD) {
+ const auto *C = LD.getSingleClause<OMPScheduleClause>();
+ if (!C)
+ return true;
+
+ switch (C->getScheduleKind()) {
+ case OMPC_SCHEDULE_guided:
+ case OMPC_SCHEDULE_runtime:
+ return false;
+ case OMPC_SCHEDULE_auto:
+ return true;
+ default: {
+ const Expr *ChunkSize = C->getChunkSize();
+ Expr::EvalResult Result;
+ if (!ChunkSize || !ChunkSize->EvaluateAsInt(Result, Ctx))
+ return false;
+ return Result.Val.getInt().getLimitedValue() == 1;
+ }
+ }
+}
+
static bool isNoLoopEligible(ASTContext &Ctx, const OMPExecutableDirective &D) {
const LangOptions &LangOpts = Ctx.getLangOpts();
if (!LangOpts.OpenMPTeamSubscription || !LangOpts.OpenMPThreadSubscription)
return false;
- OpenMPDirectiveKind DKind = D.getDirectiveKind();
- if (DKind != OMPD_target_teams_distribute_parallel_for &&
- DKind != OMPD_target_teams_distribute_parallel_for_simd)
- return false;
+ const OMPExecutableDirective *Nested = &D;
+ while (true) {
+ if (Nested->hasClausesOfKind<OMPNumTeamsClause>() ||
+ Nested->hasClausesOfKind<OMPReductionClause>())
+ return false;
+ OpenMPDirectiveKind DKind = Nested->getDirectiveKind();
+ if (DKind != OMPD_target && DKind != OMPD_target_teams &&
+ DKind != OMPD_teams)
+ break;
+ const Stmt *Body =
+ Nested->getInnermostCapturedStmt()->getCapturedStmt()->IgnoreContainers(
+ /*IgnoreCaptured=*/true);
+ Nested = Body ? dyn_cast_or_null<OMPExecutableDirective>(
+ CGOpenMPRuntime::getSingleCompoundChild(Ctx, Body))
+ : nullptr;
+ if (!Nested)
+ return false;
+ }
- const auto &LD = cast<OMPLoopDirective>(D);
+ const auto *LD = dyn_cast<OMPLoopDirective>(Nested);
// Do not filter out 'ordered' since Sema rejects it on these directives.
- if (LD.getLoopsNumber() != 1 || LD.hasClausesOfKind<OMPNumTeamsClause>() ||
- LD.hasClausesOfKind<OMPReductionClause>() ||
- LD.getSingleClause<OMPScheduleClause>() ||
- LD.getSingleClause<OMPDistScheduleClause>())
+ if (!LD || LD->getLoopsNumber() != 1 ||
+ !isNoLoopCompatibleSchedule(Ctx, *LD) ||
+ LD->getSingleClause<OMPDistScheduleClause>())
return false;
+ if (const auto *TTLD = dyn_cast<OMPTargetTeamsGenericLoopDirective>(LD))
+ return TTLD->canBeParallelFor();
+ if (const auto *TLD = dyn_cast<OMPTeamsGenericLoopDirective>(LD))
+ return TLD->canBeParallelFor();
+ if (!isOpenMPDistributeDirective(LD->getDirectiveKind()) ||
+ !isOpenMPParallelDirective(LD->getDirectiveKind()))
+ return false;
if (const auto *TTD =
- dyn_cast<OMPTargetTeamsDistributeParallelForDirective>(&LD))
+ dyn_cast<OMPTargetTeamsDistributeParallelForDirective>(LD))
return !TTD->hasCancel();
+ if (const auto *TD = dyn_cast<OMPTeamsDistributeParallelForDirective>(LD))
+ return !TD->hasCancel();
+ if (const auto *DD = dyn_cast<OMPDistributeParallelForDirective>(LD))
+ return !DD->hasCancel();
return true;
}
diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp
index 0f0802de3d19a..daf71052847c1 100644
--- a/clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -9033,30 +9033,6 @@ void CodeGenFunction::EmitOMPParallelGenericLoopDirective(
checkForLastprivateConditionalUpdate(*this, S);
}
-void CodeGenFunction::EmitOMPTeamsGenericLoopDirective(
- const OMPTeamsGenericLoopDirective &S) {
- // To be consistent with current behavior of 'target teams loop', emit
- // 'teams loop' as if its constituent constructs are 'teams' and 'distribute'.
- auto &&CodeGenDistribute = [&S](CodeGenFunction &CGF, PrePostActionTy &) {
- CGF.EmitOMPDistributeLoop(S, emitOMPLoopBodyWithStopPoint, S.getInc());
- };
-
- // Emit teams region as a standalone region.
- auto &&CodeGen = [&S, &CodeGenDistribute](CodeGenFunction &CGF,
- PrePostActionTy &Action) {
- Action.Enter(CGF);
- OMPPrivateScope PrivateScope(CGF);
- CGF.EmitOMPReductionClauseInit(S, PrivateScope);
- (void)PrivateScope.Privatize();
- CGF.CGM.getOpenMPRuntime().emitInlinedDirective(CGF, OMPD_distribute,
- CodeGenDistribute);
- CGF.EmitOMPReductionClauseFinal(S, /*ReductionKind=*/OMPD_teams);
- };
- emitCommonOMPTeamsDirective(*this, S, OMPD_distribute, CodeGen);
- emitPostUpdateForReductionClause(*this, S,
- [](CodeGenFunction &) { return nullptr; });
-}
-
#ifndef NDEBUG
static void emitTargetTeamsLoopCodegenStatus(CodeGenFunction &CGF,
std::string StatusMsg,
@@ -9076,10 +9052,9 @@ static void emitTargetTeamsLoopCodegenStatus(CodeGenFunction &CGF,
}
#endif
-static void emitTargetTeamsGenericLoopRegionAsParallel(
- CodeGenFunction &CGF, PrePostActionTy &Action,
- const OMPTargetTeamsGenericLoopDirective &S) {
- Action.Enter(CGF);
+static void
+emitTeamsGenericLoopAsDistributeParallelFor(CodeGenFunction &CGF,
+ const OMPLoopDirective &S) {
// Emit 'teams loop' as if its constituent constructs are 'distribute,
// 'parallel, and 'for'.
auto &&CodeGenDistribute = [&S](CodeGenFunction &CGF, PrePostActionTy &) {
@@ -9107,6 +9082,42 @@ static void emitTargetTeamsGenericLoopRegionAsParallel(
[](CodeGenFunction &) { return nullptr; });
}
+void CodeGenFunction::EmitOMPTeamsGenericLoopDirective(
+ const OMPTeamsGenericLoopDirective &S) {
+ if (CGM.getOpenMPRuntime().canPromoteToNoLoop()) {
+ emitTeamsGenericLoopAsDistributeParallelFor(*this, S);
+ return;
+ }
+
+ // To be consistent with current behavior of 'target teams loop', emit
+ // 'teams loop' as if its constituent constructs are 'teams' and 'distribute'.
+ auto &&CodeGenDistribute = [&S](CodeGenFunction &CGF, PrePostActionTy &) {
+ CGF.EmitOMPDistributeLoop(S, emitOMPLoopBodyWithStopPoint, S.getInc());
+ };
+
+ // Emit teams region as a standalone region.
+ auto &&CodeGen = [&S, &CodeGenDistribute](CodeGenFunction &CGF,
+ PrePostActionTy &Action) {
+ Action.Enter(CGF);
+ OMPPrivateScope PrivateScope(CGF);
+ CGF.EmitOMPReductionClauseInit(S, PrivateScope);
+ (void)PrivateScope.Privatize();
+ CGF.CGM.getOpenMPRuntime().emitInlinedDirective(CGF, OMPD_distribute,
+ CodeGenDistribute);
+ CGF.EmitOMPReductionClauseFinal(S, /*ReductionKind=*/OMPD_teams);
+ };
+ emitCommonOMPTeamsDirective(*this, S, OMPD_distribute, CodeGen);
+ emitPostUpdateForReductionClause(*this, S,
+ [](CodeGenFunction &) { return nullptr; });
+}
+
+static void emitTargetTeamsGenericLoopRegionAsParallel(
+ CodeGenFunction &CGF, PrePostActionTy &Action,
+ const OMPTargetTeamsGenericLoopDirective &S) {
+ Action.Enter(CGF);
+ emitTeamsGenericLoopAsDistributeParallelFor(CGF, S);
+}
+
static void emitTargetTeamsGenericLoopRegionAsDistribute(
CodeGenFunction &CGF, PrePostActionTy &Action,
const OMPTargetTeamsGenericLoopDirective &S) {
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index efc570a8933f1..50bffd60ef6cd 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -11770,7 +11770,8 @@ StmtResult SemaOpenMP::ActOnOpenMPTeamsGenericLoopDirective(
DSAStack->setParentTeamsRegionLoc(StartLoc);
return OMPTeamsGenericLoopDirective::Create(
- getASTContext(), StartLoc, EndLoc, NestedLoopCount, Clauses, AStmt, B);
+ getASTContext(), StartLoc, EndLoc, NestedLoopCount, Clauses, AStmt, B,
+ teamsLoopCanBeParallelFor(AStmt, SemaRef));
}
StmtResult SemaOpenMP::ActOnOpenMPTargetTeamsGenericLoopDirective(
diff --git a/clang/lib/Serialization/ASTReaderStmt.cpp b/clang/lib/Serialization/ASTReaderStmt.cpp
index 26aee27c65e82..4fe89f653e3a9 100644
--- a/clang/lib/Serialization/ASTReaderStmt.cpp
+++ b/clang/lib/Serialization/ASTReaderStmt.cpp
@@ -2957,6 +2957,7 @@ void ASTStmtReader::VisitOMPGenericLoopDirective(OMPGenericLoopDirective *D) {
void ASTStmtReader::VisitOMPTeamsGenericLoopDirective(
OMPTeamsGenericLoopDirective *D) {
VisitOMPLoopDirective(D);
+ D->setCanBeParallelFor(Record.readBool());
}
void ASTStmtReader::VisitOMPTargetTeamsGenericLoopDirective(
diff --git a/clang/lib/Serialization/ASTWriterStmt.cpp b/clang/lib/Serialization/ASTWriterStmt.cpp
index 4b4d179fb406a..876f68d46c8cb 100644
--- a/clang/lib/Serialization/ASTWriterStmt.cpp
+++ b/clang/lib/Serialization/ASTWriterStmt.cpp
@@ -3103,6 +3103,7 @@ void ASTStmtWriter::VisitOMPGenericLoopDirective(OMPGenericLoopDirective *D) {
void ASTStmtWriter::VisitOMPTeamsGenericLoopDirective(
OMPTeamsGenericLoopDirective *D) {
VisitOMPLoopDirective(D);
+ Record.writeBool(D->canBeParallelFor());
Code = serialization::STMT_OMP_TEAMS_GENERIC_LOOP_DIRECTIVE;
}
diff --git a/clang/test/OpenMP/target_no_loop.cpp b/clang/test/OpenMP/target_no_loop.cpp
index ca21a8793ab8b..5d913400f0fd1 100644
--- a/clang/test/OpenMP/target_no_loop.cpp
+++ b/clang/test/OpenMP/target_no_loop.cpp
@@ -11,6 +11,21 @@
// RUN: -emit-llvm %s -o - | FileCheck %s \
// RUN: --check-prefixes=NOLOOP
+// RUN: %clang_cc1 -fopenmp -x c++ -triple amdgcn-amd-amdhsa \
+// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp-is-target-device \
+// RUN: -fopenmp-host-ir-file-path %t-host.bc \
+// RUN: -fopenmp-assume-teams-oversubscription \
+// RUN: -fopenmp-assume-threads-oversubscription \
+// RUN: -emit-pch %s -o %t.pch
+
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple amdgcn-amd-amdhsa \
+// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp-is-target-device \
+// RUN: -fopenmp-host-ir-file-path %t-host.bc \
+// RUN: -fopenmp-assume-teams-oversubscription \
+// RUN: -fopenmp-assume-threads-oversubscription \
+// RUN: -include-pch %t.pch -emit-llvm %s -o - | FileCheck %s \
+// RUN: --check-prefixes=NOLOOP
+
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple amdgcn-amd-amdhsa \
// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -fopenmp-is-target-device \
// RUN: -fopenmp-host-ir-file-path %t-host.bc \
@@ -33,6 +48,11 @@
// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+int foo(int i);
+
void promotable_no_loop(int *array, int n) {
#pragma omp target teams distribute parallel for
for (int i = 0; i < 1024; ++i)
@@ -97,6 +117,42 @@ void promotable_no_loop(int *array, int n) {
for (i = 0; i < 1024; ++i)
array[i] = i + 1;
}
+
+#pragma omp target teams loop
+ for (int i = 0; i < 1024; ++i)
+ array[i] = i + 1;
+
+#pragma omp target
+ {
+#pragma omp teams loop
+ for (int i = 0; i < 1024; ++i)
+ array[i] = i + 1;
+ }
+
+#pragma omp target teams
+ {
+#pragma omp distribute parallel for
+ for (int i = 0; i < 1024; ++i)
+ array[i] = i + 1;
+ }
+
+#pragma omp target
+ {
+#pragma omp teams
+ {
+#pragma omp distribute parallel for
+ for (int i = 0; i < 1024; ++i)
+ array[i] = i + 1;
+ }
+ }
+
+#pragma omp target teams distribute parallel for schedule(static, 1)
+ for (int i = 0; i < 1024; ++i)
+ array[i] = i + 1;
+
+#pragma omp target teams distribute parallel for schedule(auto)
+ for (int i = 0; i < 1024; ++i)
+ array[i] = i + 1;
}
void non_promotable(int *array) {
@@ -122,108 +178,123 @@ void non_promotable(int *array) {
array[i] = i + 1;
#pragma omp cancel for
}
+
+#pragma omp target
+ {
+#pragma omp teams loop
+ for (int i = 0; i < 1024; ++i)
+ array[i] = foo(i);
+ }
}
-// NOLOOP-COUNT-10: promotable_no_loop{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 6
-// NOLOOP-COUNT-5: non_promotable{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 2
+#endif
+
+// NOLOOP-COUNT-16: promotable_no_loop{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 6
+// NOLOOP-COUNT-6: non_promotable{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 2
-// SPMD-COUNT-10: promotable_no_loop{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 2
-// SPMD-COUNT-5: non_promotable{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 2
+// SPMD-COUNT-16: promotable_no_loop{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 2
+// SPMD-COUNT-6: non_promotable{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 2
// no clause
-// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l37_{{.*}})
+// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l57_{{.*}})
// NOLOOP: omp.loop.exit:
// NOLOOP-NEXT: ret void
-// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l37_{{.*}}, i32 0, i32 0, i8 1)
+// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l57_{{.*}}, i32 0, i32 0, i8 1)
// NOLOOP: omp_loop.after:
// NOLOOP-NEXT: ret void
// simd
-// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l41_{{.*}})
+// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l61_{{.*}})
// NOLOOP: omp.loop.exit:
// NOLOOP-NEXT: store i32 1024, ptr %i
// NOLOOP-NEXT: ret void
-// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l41_{{.*}}, i32 0, i32 0, i8 1)
+// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l61_{{.*}}, i32 0, i32 0, i8 1)
// NOLOOP: omp_loop.after:
// NOLOOP-NEXT: ret void
// nowait
-// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l45_{{.*}})
+// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l65_{{.*}})
// NOLOOP: omp.loop.exit:
// NOLOOP-NEXT: ret void
-// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l45_{{.*}}, i32 0, i32 0, i8 1)
+// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l65_{{.*}}, i32 0, i32 0, i8 1)
// NOLOOP: omp_loop.exit:
// NOLOOP-NEXT: br label %omp_loop.after
// NOLOOP: omp_loop.after:
// NOLOOP-NEXT: ret void
// runtime trip count
-// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l49_{{.*}})
+// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l69_{{.*}})
// NOLOOP: omp.loop.exit:
// NOLOOP-NEXT: br label %omp.precond.end
// NOLOOP: [[LAST:%.*]] = load i32, ptr %.capture_expr.1.ascast
// NOLOOP-NEXT: [[TC:%.*]] = add nsw i32 [[LAST]], 1
-// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l49_{{.*}}, i32 [[TC]], i32 %{{.*}}, i32 0, i32 0, i8 1)
+// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l69_{{.*}}, i32 [[TC]], i32 %{{.*}}, i32 0, i32 0, i8 1)
// NOLOOP: omp_loop.after:
// NOLOOP-NEXT: ret void
// private
-// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l55_{{.*}})
+// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l75_{{.*}})
// NOLOOP: omp.loop.exit:
// NOLOOP-NEXT: ret void
// NOLOOP: store ptr %tmp1.ascast, ptr addrspace(5) %gep_tmp1.ascast
-// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l55_{{.*}}, i32 0, i32 0, i8 1)
+// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l75_{{.*}}, i32 0, i32 0, i8 1)
// NOLOOP: omp_loop.after:
// NOLOOP-NEXT: ret void
// NOLOOP: store i32 {{.*}}, ptr %loadgep_tmp1.ascast
// lambda
-// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l64_{{.*}})
+// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l84_{{.*}})
// NOLOOP: omp.loop.exit:
// NOLOOP-NEXT: ret void
// NOLOOP: [[SET:%.*]] = load ptr, ptr %set.addr.ascast
// NOLOOP-NEXT: [[FIELD:%.*]] = getelementptr inbounds nuw %class.anon, ptr [[SET]], i32 0, i32 0
// NOLOOP-NEXT: store ptr %array.addr.ascast, ptr [[FIELD]]
-// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l64_{{.*}}, i32 0, i32 0, i8 1)
+// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l84_{{.*}}, i32 0, i32 0, i8 1)
// NOLOOP: omp_loop.after:
// NOLOOP-NEXT: ret void
// lastprivate loop counter
-// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l71_{{.*}})
+// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l91_{{.*}})
// NOLOOP: omp.loop.exit:
// NOLOOP-NEXT: store i32 1024, ptr %i
// NOLOOP-NEXT: @__kmpc_free_shared(ptr %i{{.*}})
// NOLOOP-NEXT: ret void
-// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l71_{{.*}}, i32 0, i32 0, i8 1)
+// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l91_{{.*}}, i32 0, i32 0, i8 1)
// NOLOOP: @__kmpc_barrier
// NOLOOP: omp_loop.after:
// NOLOOP-NEXT: ret void
// lastprivate scalar
-// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l78_{{.*}})
+// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l98_{{.*}})
// NOLOOP: omp.loop.exit:
// NOLOOP-NEXT: @__kmpc_free_shared(ptr %last{{.*}})
// NOLOOP-NEXT: ret void
-// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l78_{{.*}}, i32 0, i32 0, i8 1)
+// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l98_{{.*}}, i32 0, i32 0, i8 1)
// NOLOOP: @__kmpc_barrier
// NOLOOP: store {{.*}}, ptr %last.
// NOLOOP-NEXT: %.omp.lastprivate.done
// lastprivate scalar, nowait
-// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l87_{{.*}})
+// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l107_{{.*}})
// NOLOOP: omp.loop.exit:
// NOLOOP-NEXT: @__kmpc_free_shared(ptr %last{{.*}})
// NOLOOP-NEXT: ret void
-// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l87_{{.*}}, i32 0, i32 0, i8 1)
+// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l107_{{.*}}, i32 0, i32 0, i8 1)
// NOLOOP: @__kmpc_barrier
// NOLOOP: store {{.*}}, ptr %last.
// NOLOOP-NEXT: %.omp.lastprivate.done
// simd linear loop counter
-// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l96_{{.*}})
+// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l116_{{.*}})
// NOLOOP: omp.loop.exit:
// NOLOOP-NEXT: store i32 1024, ptr %i
// NOLOOP-NEXT: ret void
-// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l96_{{.*}}, i32 0, i32 0, i8 1)
+// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l116_{{.*}}, i32 0, i32 0, i8 1)
// NOLOOP: omp_loop.after:
// NOLOOP-NEXT: ret void
+
+// target teams loop
+// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}promotable_no_loop{{.*}}_l121_{{.*}})
+
+// target + teams loop
+// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}promotable_no_loop{{.*}}_l125_{{.*}})
More information about the llvm-branch-commits
mailing list