https://github.com/nicebert updated https://github.com/llvm/llvm-project/pull/224650
>From e032dbe019c3b3c41e3ef9e944f38a9e4c37bcfb Mon Sep 17 00:00:00 2001 From: Nicole Aschenbrenner <[email protected]> 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 | 76 +++++++++++--- 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, 213 insertions(+), 68 deletions(-) diff --git a/clang/include/clang/AST/StmtOpenMP.h b/clang/include/clang/AST/StmtOpenMP.h index a070d97a6db1f3..55e56073b4cb21 100644 --- a/clang/include/clang/AST/StmtOpenMP.h +++ b/clang/include/clang/AST/StmtOpenMP.h @@ -6659,6 +6659,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. @@ -6680,6 +6682,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. /// @@ -6694,7 +6699,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. @@ -6708,6 +6713,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 5ba11b2d8c42fa..efc444b802afb8 100644 --- a/clang/lib/AST/StmtOpenMP.cpp +++ b/clang/lib/AST/StmtOpenMP.cpp @@ -2638,7 +2638,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, @@ -2680,6 +2680,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 97c43734717642..1dd2616f79bf0d 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp @@ -706,26 +706,61 @@ 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 const OMPLoopDirective * checkAndGetNoLoopDirective(ASTContext &Ctx, const OMPExecutableDirective &D) { const LangOptions &LangOpts = Ctx.getLangOpts(); if (!LangOpts.OpenMPTeamSubscription || !LangOpts.OpenMPThreadSubscription) return nullptr; - OpenMPDirectiveKind DKind = D.getDirectiveKind(); - if (DKind != OMPD_target_teams_distribute_parallel_for && - DKind != OMPD_target_teams_distribute_parallel_for_simd) - return nullptr; + const OMPExecutableDirective *Nested = &D; + while (true) { + if (Nested->hasClausesOfKind<OMPNumTeamsClause>() || + Nested->hasClausesOfKind<OMPReductionClause>()) + return nullptr; + 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 nullptr; + } - 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 nullptr; - for (const auto *C : LD.getClausesOfKind<OMPIfClause>()) { + for (const auto *C : LD->getClausesOfKind<OMPIfClause>()) { if (C->getNameModifier() != OMPD_unknown && C->getNameModifier() != OMPD_parallel && C->getNameModifier() != OMPD_teams) @@ -736,10 +771,25 @@ checkAndGetNoLoopDirective(ASTContext &Ctx, const OMPExecutableDirective &D) { return nullptr; } + auto getNoLoopDirective = [LD](bool Promotable) -> const OMPLoopDirective * { + return Promotable ? LD : nullptr; + }; + + if (const auto *TTLD = dyn_cast<OMPTargetTeamsGenericLoopDirective>(LD)) + return getNoLoopDirective(TTLD->canBeParallelFor()); + if (const auto *TLD = dyn_cast<OMPTeamsGenericLoopDirective>(LD)) + return getNoLoopDirective(TLD->canBeParallelFor()); + if (!isOpenMPDistributeDirective(LD->getDirectiveKind()) || + !isOpenMPParallelDirective(LD->getDirectiveKind())) + return nullptr; if (const auto *TTD = - dyn_cast<OMPTargetTeamsDistributeParallelForDirective>(&LD)) - return TTD->hasCancel() ? nullptr : &LD; - return &LD; + dyn_cast<OMPTargetTeamsDistributeParallelForDirective>(LD)) + return getNoLoopDirective(!TTD->hasCancel()); + if (const auto *TD = dyn_cast<OMPTeamsDistributeParallelForDirective>(LD)) + return getNoLoopDirective(!TD->hasCancel()); + if (const auto *DD = dyn_cast<OMPDistributeParallelForDirective>(LD)) + return getNoLoopDirective(!DD->hasCancel()); + return LD; } void CGOpenMPRuntimeGPU::emitNonSPMDKernel(const OMPExecutableDirective &D, diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp index 4cd348bbb5f6ad..e654e711db4517 100644 --- a/clang/lib/CodeGen/CGStmtOpenMP.cpp +++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp @@ -9054,30 +9054,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, @@ -9097,10 +9073,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 &) { @@ -9128,6 +9103,42 @@ static void emitTargetTeamsGenericLoopRegionAsParallel( [](CodeGenFunction &) { return nullptr; }); } +void CodeGenFunction::EmitOMPTeamsGenericLoopDirective( + const OMPTeamsGenericLoopDirective &S) { + if (CGM.getOpenMPRuntime().canPromoteToNoLoop(S)) { + 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 efe8c9338b5a9e..5ff9ab02f879c3 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 26aee27c65e82f..4fe89f653e3a9f 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 4b4d179fb406ac..876f68d46c8cb3 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 cb8953bc3588c6..1001547a47baf1 100644 --- a/clang/test/OpenMP/target_no_loop.cpp +++ b/clang/test/OpenMP/target_no_loop.cpp @@ -21,6 +21,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 \ @@ -43,6 +58,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) @@ -111,6 +131,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, int n) { @@ -144,6 +200,13 @@ void non_promotable(int *array, int n) { #pragma omp target teams distribute parallel for if(n > 0) 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] = foo(i); + } } #pragma omp declare target @@ -161,13 +224,15 @@ void promotable_after_reduction(int *array) { array[i] = i + offset; } -// NOLOOP-COUNT-11: promotable_no_loop{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 6 -// NOLOOP-COUNT-7: non_promotable{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 2 +#endif + +// NOLOOP-COUNT-17: promotable_no_loop{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 6 +// NOLOOP-COUNT-8: non_promotable{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 2 // NOLOOP: promotable_after_reduction{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 2 // NOLOOP: promotable_after_reduction{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 6 -// SPMD-COUNT-11: promotable_no_loop{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 2 -// SPMD-COUNT-7: non_promotable{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 2 +// SPMD-COUNT-17: promotable_no_loop{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 2 +// SPMD-COUNT-8: non_promotable{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 2 // SPMD-COUNT-2: promotable_after_reduction{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 2 // reduction before a promoted kernel @@ -187,99 +252,105 @@ void promotable_after_reduction(int *array) { // HOST-NOFLAGS-LABEL: define {{.*}}@_Z18promotable_no_loopPii( // no clause -// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l47_{{.*}}) +// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l67_{{.*}}) // NOLOOP: omp.loop.exit: // NOLOOP-NEXT: ret void -// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l47_{{.*}}, i32 0, i32 0, i8 1) +// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l67_{{.*}}, i32 0, i32 0, i8 1) // NOLOOP: omp_loop.after: // NOLOOP-NEXT: ret void // simd -// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l51_{{.*}}) +// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l71_{{.*}}) // NOLOOP: omp.loop.exit: // NOLOOP-NEXT: store i32 1024, ptr %i // NOLOOP-NEXT: ret void -// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l51_{{.*}}, i32 0, i32 0, i8 1) +// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l71_{{.*}}, i32 0, i32 0, i8 1) // NOLOOP: omp_loop.after: // NOLOOP-NEXT: ret void // nowait -// 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: @__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.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{{.*}}_l59_{{.*}}) +// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l79_{{.*}}) // 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({{.*}}_l59_{{.*}}, i32 [[TC]], i32 %{{.*}}, i32 0, i32 0, i8 1) +// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l79_{{.*}}, 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{{.*}}_l65_{{.*}}) +// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l85_{{.*}}) // 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({{.*}}_l65_{{.*}}, i32 0, i32 0, i8 1) +// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l85_{{.*}}, 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{{.*}}_l74_{{.*}}) +// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l94_{{.*}}) // 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({{.*}}_l74_{{.*}}, i32 0, i32 0, i8 1) +// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l94_{{.*}}, 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{{.*}}_l85_{{.*}}) +// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l105_{{.*}}) // 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({{.*}}_l85_{{.*}}, i32 0, i32 0, i8 1) +// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l105_{{.*}}, 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{{.*}}_l92_{{.*}}) +// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l112_{{.*}}) // NOLOOP: omp.loop.exit: // NOLOOP-NEXT: @__kmpc_free_shared(ptr %last{{.*}}) // NOLOOP-NEXT: ret void -// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l92_{{.*}}, i32 0, i32 0, i8 1) +// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l112_{{.*}}, 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{{.*}}_l101_{{.*}}) +// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l121_{{.*}}) // NOLOOP: omp.loop.exit: // NOLOOP-NEXT: @__kmpc_free_shared(ptr %last{{.*}}) // NOLOOP-NEXT: ret void -// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l101_{{.*}}, i32 0, i32 0, i8 1) +// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l121_{{.*}}, 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{{.*}}_l110_{{.*}}) +// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l130_{{.*}}) // NOLOOP: omp.loop.exit: // NOLOOP-NEXT: store i32 1024, ptr %i // NOLOOP-NEXT: ret void -// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l110_{{.*}}, i32 0, i32 0, i8 1) +// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l130_{{.*}}, 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{{.*}}_l135_{{.*}}) + +// target + teams loop +// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}promotable_no_loop{{.*}}_l139_{{.*}}) _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
