[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
Mon Oct 5 23:57:15 PDT 2026
https://github.com/nicebert updated https://github.com/llvm/llvm-project/pull/224650
>From a6562c24e3365286a922492365b47546873601de 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/CGOpenMPRuntime.h | 9 +-
clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp | 94 +++++++++++++---
clang/lib/CodeGen/CGOpenMPRuntimeGPU.h | 14 ++-
clang/lib/CodeGen/CGStmtOpenMP.cpp | 80 ++++++-------
clang/lib/Sema/SemaOpenMP.cpp | 3 +-
clang/lib/Serialization/ASTReaderStmt.cpp | 1 +
clang/lib/Serialization/ASTWriterStmt.cpp | 1 +
clang/test/OpenMP/target_no_loop.c | 130 +++++++++++++++++++++-
10 files changed, 278 insertions(+), 68 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/CGOpenMPRuntime.h b/clang/lib/CodeGen/CGOpenMPRuntime.h
index 9a93f5c5c2e82..29cc5c0bc2711 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.h
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.h
@@ -669,11 +669,10 @@ class CGOpenMPRuntime {
return false;
};
- /// Check whether a target kernel can be promoted to a "no-loop" SPMD kernel,
- /// mirroring Flang's MLIR promotion path.
- virtual bool canPromoteToNoLoop(const OMPExecutableDirective &D) const {
- return false;
- }
+ /// Check whether the target kernel being emitted is tagged SPMD_NO_LOOP, to
+ /// complete the promotion to a "no-loop" SPMD kernel, mirroring Flang's MLIR
+ /// promotion path.
+ virtual bool canPromoteToNoLoop() const { return false; }
/// Get call to __kmpc_alloc_shared
virtual std::pair<llvm::Value *, llvm::Value *>
diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
index fa6fcdab47b18..36bbaad10efe9 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
@@ -706,6 +706,75 @@ 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;
+
+ 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 = dyn_cast<OMPLoopDirective>(Nested);
+ 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))
+ 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;
+}
+
void CGOpenMPRuntimeGPU::emitNonSPMDKernel(const OMPExecutableDirective &D,
StringRef ParentName,
llvm::Function *&OutlinedFn,
@@ -745,6 +814,7 @@ void CGOpenMPRuntimeGPU::emitNonSPMDKernel(const OMPExecutableDirective &D,
emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID,
IsOffloadEntry, CodeGen);
IsInTTDRegion = false;
+ KernelAttrs = {};
}
void CGOpenMPRuntimeGPU::emitBareKernelEnvironment(
@@ -762,19 +832,19 @@ void CGOpenMPRuntimeGPU::emitBareKernelEnvironment(
void CGOpenMPRuntimeGPU::emitKernelInit(const OMPExecutableDirective &D,
CodeGenFunction &CGF,
EntryFunctionState &EST, bool IsSPMD) {
- llvm::OpenMPIRBuilder::TargetKernelDefaultAttrs Attrs;
- if (IsSPMD && canPromoteToNoLoop(D))
- Attrs.ExecFlags =
+ KernelAttrs = {};
+ if (IsSPMD && isNoLoopEligible(CGM.getContext(), D))
+ KernelAttrs.ExecFlags =
llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_SPMD_NO_LOOP;
else
- Attrs.ExecFlags =
+ KernelAttrs.ExecFlags =
IsSPMD ? llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_SPMD
: llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_GENERIC;
- computeMinAndMaxThreadsAndTeams(D, CGF, Attrs);
+ computeMinAndMaxThreadsAndTeams(D, CGF, KernelAttrs);
CGBuilderTy &Bld = CGF.Builder;
- Bld.restoreIP(OMPBuilder.createTargetInit(Bld, Attrs));
+ Bld.restoreIP(OMPBuilder.createTargetInit(Bld, KernelAttrs));
if (!IsSPMD)
emitGenericVarsProlog(CGF, EST.Loc);
}
@@ -863,6 +933,7 @@ void CGOpenMPRuntimeGPU::emitSPMDKernel(const OMPExecutableDirective &D,
emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID,
IsOffloadEntry, CodeGen);
IsInTTDRegion = false;
+ KernelAttrs = {};
}
void CGOpenMPRuntimeGPU::emitTargetOutlinedFunction(
@@ -1163,17 +1234,6 @@ bool CGOpenMPRuntimeGPU::isDelayedVariableLengthDecl(CodeGenFunction &CGF,
return llvm::is_contained(I->getSecond().DelayedVariableLengthDecls, VD);
}
-bool CGOpenMPRuntimeGPU::canPromoteToNoLoop(
- const OMPExecutableDirective &D) const {
- OpenMPDirectiveKind DKind = D.getDirectiveKind();
- const LangOptions &LangOpts = CGM.getLangOpts();
- return (DKind == OMPD_target_teams_distribute_parallel_for ||
- DKind == OMPD_target_teams_distribute_parallel_for_simd) &&
- LangOpts.OpenMPTeamSubscription && LangOpts.OpenMPThreadSubscription &&
- !D.hasClausesOfKind<OMPNumTeamsClause>() &&
- !D.hasClausesOfKind<OMPReductionClause>();
-}
-
std::pair<llvm::Value *, llvm::Value *>
CGOpenMPRuntimeGPU::getKmpcAllocShared(CodeGenFunction &CGF,
const VarDecl *VD) {
diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.h b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.h
index d40d15767552e..669fc34ceecb1 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.h
+++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.h
@@ -146,9 +146,13 @@ class CGOpenMPRuntimeGPU : public CGOpenMPRuntime {
bool isDelayedVariableLengthDecl(CodeGenFunction &CGF,
const VarDecl *VD) const override;
- /// Check whether a target kernel can be promoted to a "no-loop" SPMD kernel,
- /// mirroring Flang's MLIR promotion path.
- bool canPromoteToNoLoop(const OMPExecutableDirective &D) const override;
+ /// Check whether the target kernel being emitted is tagged SPMD_NO_LOOP, to
+ /// complete the promotion to a "no-loop" SPMD kernel, mirroring Flang's MLIR
+ /// promotion path.
+ bool canPromoteToNoLoop() const override {
+ return KernelAttrs.ExecFlags ==
+ llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_SPMD_NO_LOOP;
+ }
/// Get call to __kmpc_alloc_shared
std::pair<llvm::Value *, llvm::Value *>
@@ -382,6 +386,10 @@ class CGOpenMPRuntimeGPU : public CGOpenMPRuntime {
/// - otherwise.
bool IsInTTDRegion = false;
+ /// The default attributes of the target region being emitted, filled in by
+ /// the kernel prologue and read back while the region's body is emitted.
+ llvm::OpenMPIRBuilder::TargetKernelDefaultAttrs KernelAttrs;
+
/// Map between an outlined function and its wrapper.
llvm::DenseMap<llvm::Function *, llvm::Function *> WrapperFunctionsMap;
diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp
index 9f740b98001ce..bfba75e919d6b 100644
--- a/clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -98,15 +98,6 @@ hasLoopLastprivateVariable(const OMPLoopDirective &S,
return false;
}
-static bool canEmitGPUNoLoopKernel(CodeGenModule &CGM,
- const OMPLoopDirective &S) {
- const auto *D = dyn_cast<OMPTargetTeamsDistributeParallelForDirective>(&S);
- return S.getLoopsNumber() == 1 &&
- CGM.getOpenMPRuntime().canPromoteToNoLoop(S) &&
- !S.getSingleClause<OMPScheduleClause>() &&
- !S.getSingleClause<OMPDistScheduleClause>() && !(D && D->hasCancel());
-}
-
namespace {
/// Lexical scope for OpenMP executable constructs, that handles correct codegen
/// for captured expressions.
@@ -3995,7 +3986,7 @@ emitInnerParallelForWhenCombined(CodeGenFunction &CGF,
CodeGenFunction::OMPCancelStackRAII CancelRegion(CGF, EKind, HasCancel);
CodeGenModule &CGM = CGF.CGM;
- if (canEmitGPUNoLoopKernel(CGM, S)) {
+ if (CGM.getOpenMPRuntime().canPromoteToNoLoop()) {
// Prepare the loop variables and their privatization.
emitLoopIterationspaceVars(CGF, S);
OMPLoopScope PreInitScope(CGF, S);
@@ -6647,7 +6638,7 @@ void CodeGenFunction::EmitOMPDistributeLoop(const OMPLoopDirective &S,
EmitOMPReductionClauseInit(S, LoopScope);
// Defer lastprivate init to the parallel region for no-loop
- const bool NoLoopKernel = canEmitGPUNoLoopKernel(CGM, S);
+ const bool NoLoopKernel = CGM.getOpenMPRuntime().canPromoteToNoLoop();
if (!NoLoopKernel)
HasLastprivateClause = EmitOMPLastprivateClauseInit(S, LoopScope);
@@ -9042,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,
@@ -9085,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 &) {
@@ -9116,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.c b/clang/test/OpenMP/target_no_loop.c
index 41b7ab96e8867..223a7b0631a23 100644
--- a/clang/test/OpenMP/target_no_loop.c
+++ b/clang/test/OpenMP/target_no_loop.c
@@ -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 no_loop(int *array) {
#pragma omp target teams distribute parallel for
for (int i = 0; i < 1024; ++i)
@@ -76,9 +96,114 @@ void no_loop_lastprivate_scalar_nowait(int *array) {
}
}
+void fused_teams_loop(int *array) {
+#pragma omp target teams loop
+ for (int i = 0; i < 1024; ++i)
+ array[i] = i + 1;
+}
+
+void split_teams_loop(int *array) {
+#pragma omp target
+ {
+#pragma omp teams loop
+ for (int i = 0; i < 1024; ++i)
+ array[i] = i + 1;
+ }
+}
+
+void split_teams_loop_call(int *array) {
+#pragma omp target
+ {
+#pragma omp teams loop
+ for (int i = 0; i < 1024; ++i)
+ array[i] = foo(i);
+ }
+}
+
+void split_teams_dpf(int *array) {
+#pragma omp target teams
+ {
+#pragma omp distribute parallel for
+ for (int i = 0; i < 1024; ++i)
+ array[i] = i + 1;
+ }
+}
+
+void split_target_teams_dpf(int *array) {
+#pragma omp target
+ {
+#pragma omp teams
+ {
+#pragma omp distribute parallel for
+ for (int i = 0; i < 1024; ++i)
+ array[i] = i + 1;
+ }
+ }
+}
+
+void spmd_collapse2(int *array) {
+#pragma omp target teams distribute parallel for collapse(2)
+ for (int i = 0; i < 32; ++i)
+ for (int j = 0; j < 32; ++j)
+ array[i * 32 + j] = i + j;
+}
+
+void spmd_schedule_static(int *array) {
+#pragma omp target teams distribute parallel for schedule(static)
+ for (int i = 0; i < 1024; ++i)
+ array[i] = i + 1;
+}
+
+void no_loop_schedule_static_one(int *array) {
+#pragma omp target teams distribute parallel for schedule(static, 1)
+ for (int i = 0; i < 1024; ++i)
+ array[i] = i + 1;
+}
+
+void no_loop_schedule_auto(int *array) {
+#pragma omp target teams distribute parallel for schedule(auto)
+ for (int i = 0; i < 1024; ++i)
+ array[i] = i + 1;
+}
+
+void spmd_schedule_guided(int *array) {
+#pragma omp target teams distribute parallel for schedule(guided)
+ for (int i = 0; i < 1024; ++i)
+ array[i] = i + 1;
+}
+
+void spmd_dist_schedule(int *array) {
+#pragma omp target teams distribute parallel for dist_schedule(static)
+ for (int i = 0; i < 1024; ++i)
+ array[i] = i + 1;
+}
+
+void spmd_cancel(int *array) {
+#pragma omp target teams distribute parallel for
+ for (int i = 0; i < 1024; ++i) {
+ array[i] = i + 1;
+#pragma omp cancel for
+ }
+}
+
+#endif
+
// NOLOOP: no_loop_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 6
// NOLOOP: no_loop_simd_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 6
+// NOLOOP: fused_teams_loop_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 6
+// NOLOOP: split_teams_loop_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 6
+// NOLOOP: split_teams_loop_call_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 2
+// NOLOOP: split_teams_dpf_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 6
+// NOLOOP: split_target_teams_dpf_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 6
+// NOLOOP: spmd_collapse2_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 2
+// NOLOOP: spmd_schedule_static_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 2
+// NOLOOP: no_loop_schedule_static_one_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 6
+// NOLOOP: no_loop_schedule_auto_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 6
+// NOLOOP: spmd_schedule_guided_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 2
+// NOLOOP: spmd_dist_schedule_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 2
+// NOLOOP: spmd_cancel_l{{[0-9]+}}_kernel_environment {{.*}} i8 0, i8 1, i8 2
+
// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}no_loop_l{{[0-9]+}}{{.*}})
// NOLOOP: omp.loop.exit:
// NOLOOP-NEXT: ret void
@@ -131,4 +256,7 @@ void no_loop_lastprivate_scalar_nowait(int *array) {
// NOLOOP: store {{.*}}, ptr %last.
// NOLOOP-NEXT: %.omp.lastprivate.done
-// SPMD-COUNT-6: _kernel_environment {{.*}} i8 0, i8 1, i8 2
+// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}fused_teams_loop_l{{[0-9]+}}{{.*}})
+// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}split_teams_loop_l{{[0-9]+}}{{.*}})
+
+// SPMD-COUNT-18: _kernel_environment {{.*}} i8 0, i8 1, i8 2
More information about the llvm-branch-commits
mailing list