[clang] [llvm] [clang][OpenMP] Add no-loop SPMD kernel promotion (PR #224041)
Nicole Aschenbrenner via llvm-commits
llvm-commits at lists.llvm.org
Fri Sep 25 02:29:20 PDT 2026
https://github.com/nicebert updated https://github.com/llvm/llvm-project/pull/224041
>From 34338f853176b085795e59098a7300e1d5aea1cc Mon Sep 17 00:00:00 2001
From: Nicole Aschenbrenner <nicole.aschenbrenner at amd.com>
Date: Tue, 15 Sep 2026 10:34:44 -0500
Subject: [PATCH 1/2] [clang][OpenMP] Add no-loop SPMD kernel promotion
A target teams distribute parallel for that is guaranteed a thread for
every iteration does not need the loop around its body. Flang already
drops it and runs the region as a no-loop kernel.
Enable the same optimization for Clang through mirroring Flang's MLIR
promotion using OpenMPIRBuilder. The kernel is tagged SPMD_NO_LOOP, so
the runtime sizes the grid to the iteration space, and the body is
emitted without a loop around it. The canonical loop it consumes is
reconstructed in the no-loop branch rather than taken from an
OMPCanonicalLoop node, so the promotion does not require
-fopenmp-enable-irbuilder.
Restrict offload entry creation to module level finalize, preventing
asserts on missing offload entries from nested CodeGenFunction
finalizing before module completion.
---
clang/lib/CodeGen/CGOpenMPRuntime.h | 6 +
clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp | 24 ++-
clang/lib/CodeGen/CGOpenMPRuntimeGPU.h | 4 +
clang/lib/CodeGen/CGStmtOpenMP.cpp | 172 +++++++++++++-----
clang/lib/CodeGen/CodeGenFunction.cpp | 18 +-
clang/lib/CodeGen/CodeGenFunction.h | 3 +-
clang/test/OpenMP/target_no_loop.c | 81 +++++++++
.../llvm/Frontend/OpenMP/OMPIRBuilder.h | 8 +
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 2 +-
offload/test/offloading/target-no-loop.c | 91 +++++++++
10 files changed, 352 insertions(+), 57 deletions(-)
create mode 100644 clang/test/OpenMP/target_no_loop.c
create mode 100644 offload/test/offloading/target-no-loop.c
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.h b/clang/lib/CodeGen/CGOpenMPRuntime.h
index 73bbcc05f515a5..fb78f90e33878c 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.h
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.h
@@ -669,6 +669,12 @@ 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;
+ }
+
/// Get call to __kmpc_alloc_shared
virtual std::pair<llvm::Value *, llvm::Value *>
getKmpcAllocShared(CodeGenFunction &CGF, const VarDecl *VD) {
diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
index e785ee635af707..264d78b0f52452 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
@@ -751,9 +751,14 @@ void CGOpenMPRuntimeGPU::emitKernelInit(const OMPExecutableDirective &D,
CodeGenFunction &CGF,
EntryFunctionState &EST, bool IsSPMD) {
llvm::OpenMPIRBuilder::TargetKernelDefaultAttrs Attrs;
- Attrs.ExecFlags =
- IsSPMD ? llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_SPMD
- : llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_GENERIC;
+ if (IsSPMD && canPromoteToNoLoop(D))
+ Attrs.ExecFlags =
+ llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_SPMD_NO_LOOP;
+ else
+ Attrs.ExecFlags =
+ IsSPMD ? llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_SPMD
+ : llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_GENERIC;
+
computeMinAndMaxThreadsAndTeams(D, CGF, Attrs);
CGBuilderTy &Bld = CGF.Builder;
@@ -1145,6 +1150,19 @@ 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>() &&
+ !D.hasClausesOfKind<OMPLastprivateClause>() &&
+ !D.hasClausesOfKind<OMPLinearClause>();
+}
+
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 3a7ee5456a9d2e..d1f682780b9bdc 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.h
+++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.h
@@ -141,6 +141,10 @@ 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;
+
/// Get call to __kmpc_alloc_shared
std::pair<llvm::Value *, llvm::Value *>
getKmpcAllocShared(CodeGenFunction &CGF, const VarDecl *VD) override;
diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp
index 7c3b30c6cedc0e..2836911518398e 100644
--- a/clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -66,6 +66,15 @@ static bool canEmitGPUFusedDistSchedule(const CodeGenModule &CGM,
!S.getSingleClause<OMPOrderedClause>();
}
+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.
@@ -2371,6 +2380,24 @@ emitCapturedStmtCall(CodeGenFunction &ParentCGF, EmittedClosureTy Cap,
return ParentCGF.Builder.CreateCall(Cap.first, EffectiveArgs);
}
+static llvm::CanonicalLoopInfo *
+createCanonicalLoop(CodeGenFunction &CGF, llvm::Value *TripCount,
+ llvm::function_ref<void(llvm::Value *)> BodyGen) {
+ auto BodyGenCB = [&](llvm::OpenMPIRBuilder::InsertPointTy CodeGenIP,
+ llvm::Value *IndVar) {
+ CGF.Builder.restoreIP(CodeGenIP);
+ BodyGen(IndVar);
+ return llvm::Error::success();
+ };
+
+ llvm::OpenMPIRBuilder &OMPBuilder =
+ CGF.CGM.getOpenMPRuntime().getOMPBuilder();
+ llvm::CanonicalLoopInfo *CL = cantFail(
+ OMPBuilder.createCanonicalLoop(CGF.Builder, BodyGenCB, TripCount));
+ CGF.Builder.restoreIP(CL->getAfterIP());
+ return CL;
+}
+
llvm::CanonicalLoopInfo *
CodeGenFunction::EmitOMPCollapsedCanonicalLoopNest(const Stmt *S, int Depth) {
assert(Depth == 1 && "Nested loops with OpenMPIRBuilder not yet implemented");
@@ -2446,11 +2473,7 @@ void CodeGenFunction::EmitOMPCanonicalLoop(const OMPCanonicalLoop *S) {
llvm::Value *DistVal = Builder.CreateLoad(CountAddr, ".count");
// Emit the loop structure.
- llvm::OpenMPIRBuilder &OMPBuilder = CGM.getOpenMPRuntime().getOMPBuilder();
- auto BodyGen = [&, this](llvm::OpenMPIRBuilder::InsertPointTy CodeGenIP,
- llvm::Value *IndVar) {
- Builder.restoreIP(CodeGenIP);
-
+ auto BodyGen = [&, this](llvm::Value *IndVar) {
// Emit the loop body: Convert the logical iteration number to the loop
// variable and emit the body.
const DeclRefExpr *LoopVarRef = S->getLoopVarRef();
@@ -2461,14 +2484,11 @@ void CodeGenFunction::EmitOMPCanonicalLoop(const OMPCanonicalLoop *S) {
RunCleanupsScope BodyScope(*this);
EmitStmt(BodyStmt);
- return llvm::Error::success();
};
- llvm::CanonicalLoopInfo *CL =
- cantFail(OMPBuilder.createCanonicalLoop(Builder, BodyGen, DistVal));
+ llvm::CanonicalLoopInfo *CL = createCanonicalLoop(*this, DistVal, BodyGen);
// Finish up the loop.
- Builder.restoreIP(CL->getAfterIP());
ForScope.ForceCleanup();
// Remember the CanonicalLoopInfo for parent AST nodes consuming it.
@@ -2648,12 +2668,26 @@ static void emitAlignedClause(CodeGenFunction &CGF,
}
void CodeGenFunction::EmitOMPPrivateLoopCounters(
- const OMPLoopDirective &S, CodeGenFunction::OMPPrivateScope &LoopScope) {
+ const OMPLoopDirective &S, CodeGenFunction::OMPPrivateScope &LoopScope,
+ bool OnlyUnresolved) {
if (!HaveInsertPoint())
return;
auto I = S.private_counters().begin();
for (const Expr *E : S.counters()) {
- const auto *VD = cast<VarDecl>(cast<DeclRefExpr>(E)->getDecl());
+ const auto *DRE = cast<DeclRefExpr>(E);
+ const auto *VD = cast<VarDecl>(DRE->getDecl());
+ // Skip counters that already resolve, mirroring EmitDeclRefLValue's
+ // handling for these cases.
+ if (OnlyUnresolved) {
+ const VarDecl *Canonical = VD->getCanonicalDecl();
+ if (!DRE->refersToEnclosingVariableOrCapture() || !CapturedStmtInfo ||
+ LocalDeclMap.count(Canonical) ||
+ CapturedStmtInfo->lookup(Canonical)) {
+ ++I;
+ continue;
+ }
+ }
+
const auto *PrivateVD = cast<VarDecl>(cast<DeclRefExpr>(*I)->getDecl());
// Emit var without initialization.
AutoVarEmission VarEmission = EmitAutoVarAlloca(*PrivateVD);
@@ -3691,6 +3725,22 @@ static void emitDistributeParallelForDistributeInnerBoundParams(
CapturedVars.push_back(UBCast);
}
+static void emitLoopIterationspaceVars(CodeGenFunction &CGF,
+ const OMPLoopDirective &S) {
+ // Emit the loop iteration variable.
+ const auto *IVExpr = cast<DeclRefExpr>(S.getIterationVariable());
+ CGF.EmitVarDecl(*cast<VarDecl>(IVExpr->getDecl()));
+
+ // Emit the iterations count variable.
+ // If it is not a variable, Sema decided to calculate iterations count on each
+ // iteration (e.g., it is foldable into a constant).
+ if (const auto *LIExpr = dyn_cast<DeclRefExpr>(S.getLastIteration())) {
+ CGF.EmitVarDecl(*cast<VarDecl>(LIExpr->getDecl()));
+ // Emit calculation of the iterations count.
+ CGF.EmitIgnoredExpr(S.getCalcLastIteration());
+ }
+}
+
static void
emitInnerParallelForWhenCombined(CodeGenFunction &CGF,
const OMPLoopDirective &S,
@@ -3710,6 +3760,55 @@ emitInnerParallelForWhenCombined(CodeGenFunction &CGF,
HasCancel = D->hasCancel();
}
CodeGenFunction::OMPCancelStackRAII CancelRegion(CGF, EKind, HasCancel);
+
+ CodeGenModule &CGM = CGF.CGM;
+ if (canEmitGPUNoLoopKernel(CGM, S)) {
+ // Prepare the loop variables and their privatization.
+ emitLoopIterationspaceVars(CGF, S);
+ OMPLoopScope PreInitScope(CGF, S);
+
+ CodeGenFunction::OMPPrivateScope PrivateScope(CGF);
+ CGF.EmitOMPPrivateClause(S, PrivateScope);
+ CGF.EmitOMPPrivateLoopCounters(S, PrivateScope);
+ (void)PrivateScope.Privatize();
+
+ if (isOpenMPTargetExecutionDirective(EKind))
+ CGM.getOpenMPRuntime().adjustTargetSpecificDataForLambdas(CGF, S);
+
+ // Rebuild what the OMPCanonicalLoop node supplies under the IRBuilder
+ // flag: iteration count, loop, index-to-variable mapping, body.
+ const Expr *NumIterations = S.getNumIterations();
+ llvm::Value *TripCount = CGF.EmitScalarConversion(
+ CGF.EmitScalarExpr(NumIterations), NumIterations->getType(),
+ S.getIterationVariable()->getType(), S.getBeginLoc());
+ llvm::CanonicalLoopInfo *CLI =
+ createCanonicalLoop(CGF, TripCount, [&CGF, &S](llvm::Value *IndVar) {
+ llvm::BasicBlock *BodyExit = llvm::splitBBWithSuffix(
+ CGF.Builder, /*CreateBranch=*/false, ".cont");
+ CGF.EmitStoreOfScalar(IndVar,
+ CGF.EmitLValue(S.getIterationVariable()));
+ emitOMPLoopBodyWithStopPoint(CGF, S, CodeGenFunction::JumpDest());
+ CGF.Builder.CreateBr(BodyExit);
+ });
+
+ llvm::OpenMPIRBuilder::InsertPointTy AllocaIP(
+ CGF.AllocaInsertPt->getParent(), CGF.AllocaInsertPt->getIterator());
+ llvm::OpenMPIRBuilder &OMPBuilder =
+ CGM.getOpenMPRuntime().getOMPBuilder();
+
+ cantFail(OMPBuilder.applyWorkshareLoop(
+ CGF.Builder.getCurrentDebugLocation(), CLI, AllocaIP,
+ /*NeedsBarrier=*/!S.getSingleClause<OMPNowaitClause>(),
+ llvm::omp::OMP_SCHEDULE_Default,
+ /*ChunkSize=*/nullptr, /*HasSimdModifier=*/false,
+ /*HasMonotonicModifier=*/false, /*HasNonmonotonicModifier=*/false,
+ /*HasOrderedClause=*/false,
+ llvm::omp::WorksharingLoopType::DistributeForStaticLoop,
+ /*NoLoop=*/true, /*HasDistSchedule=*/false,
+ /*DistScheduleChunkSize=*/nullptr));
+ return;
+ }
+
CGF.EmitOMPWorksharingLoop(S, S.getPrevEnsureUpperBound(),
emitDistributeParallelForInnerBounds,
emitDistributeParallelForDispatchBounds);
@@ -3788,19 +3887,7 @@ bool CodeGenFunction::EmitOMPWorksharingLoop(
const OMPLoopDirective &S, Expr *EUB,
const CodeGenLoopBoundsTy &CodeGenLoopBounds,
const CodeGenDispatchBoundsTy &CGDispatchBounds) {
- // Emit the loop iteration variable.
- const auto *IVExpr = cast<DeclRefExpr>(S.getIterationVariable());
- const auto *IVDecl = cast<VarDecl>(IVExpr->getDecl());
- EmitVarDecl(*IVDecl);
-
- // Emit the iterations count variable.
- // If it is not a variable, Sema decided to calculate iterations count on each
- // iteration (e.g., it is foldable into a constant).
- if (const auto *LIExpr = dyn_cast<DeclRefExpr>(S.getLastIteration())) {
- EmitVarDecl(*cast<VarDecl>(LIExpr->getDecl()));
- // Emit calculation of the iterations count.
- EmitIgnoredExpr(S.getCalcLastIteration());
- }
+ emitLoopIterationspaceVars(*this, S);
CGOpenMPRuntime &RT = CGM.getOpenMPRuntime();
@@ -3895,6 +3982,7 @@ bool CodeGenFunction::EmitOMPWorksharingLoop(
HasChunkSizeOne = (EvaluatedChunk.getLimitedValue() == 1);
}
}
+ const auto *IVExpr = cast<DeclRefExpr>(S.getIterationVariable());
const unsigned IVSize = getContext().getTypeSize(IVExpr->getType());
const bool IVSigned = IVExpr->getType()->hasSignedIntegerRepresentation();
// OpenMP 4.5, 2.7.1 Loop Construct, Description.
@@ -6216,19 +6304,7 @@ void CodeGenFunction::EmitOMPScanDirective(const OMPScanDirective &S) {
void CodeGenFunction::EmitOMPDistributeLoop(const OMPLoopDirective &S,
const CodeGenLoopTy &CodeGenLoop,
Expr *IncExpr) {
- // Emit the loop iteration variable.
- const auto *IVExpr = cast<DeclRefExpr>(S.getIterationVariable());
- const auto *IVDecl = cast<VarDecl>(IVExpr->getDecl());
- EmitVarDecl(*IVDecl);
-
- // Emit the iterations count variable.
- // If it is not a variable, Sema decided to calculate iterations count on each
- // iteration (e.g., it is foldable into a constant).
- if (const auto *LIExpr = dyn_cast<DeclRefExpr>(S.getLastIteration())) {
- EmitVarDecl(*cast<VarDecl>(LIExpr->getDecl()));
- // Emit calculation of the iterations count.
- EmitIgnoredExpr(S.getCalcLastIteration());
- }
+ emitLoopIterationspaceVars(*this, S);
CGOpenMPRuntime &RT = CGM.getOpenMPRuntime();
@@ -6287,6 +6363,8 @@ void CodeGenFunction::EmitOMPDistributeLoop(const OMPLoopDirective &S,
!isOpenMPParallelDirective(S.getDirectiveKind()) &&
!isOpenMPTeamsDirective(S.getDirectiveKind()))
EmitOMPReductionClauseInit(S, LoopScope);
+
+ const bool NoLoopKernel = canEmitGPUNoLoopKernel(CGM, S);
HasLastprivateClause = EmitOMPLastprivateClauseInit(S, LoopScope);
EmitOMPPrivateLoopCounters(S, LoopScope);
(void)LoopScope.Privatize();
@@ -6309,12 +6387,15 @@ void CodeGenFunction::EmitOMPDistributeLoop(const OMPLoopDirective &S,
CGM.getOpenMPRuntime().getDefaultDistScheduleAndChunk(
*this, S, ScheduleKind, Chunk);
}
+ const auto *IVExpr = cast<DeclRefExpr>(S.getIterationVariable());
const unsigned IVSize = getContext().getTypeSize(IVExpr->getType());
const bool IVSigned = IVExpr->getType()->hasSignedIntegerRepresentation();
- // GPU fused schedule: omit the outer distribute loop and let the inner
- // worksharing loop schedule the flattened team/thread iteration space.
- if (canEmitGPUFusedDistSchedule(CGM, S, S.getDirectiveKind())) {
+ // omit the outer distribute loop and let the inner worksharing loop
+ // schedule the flattened team/thread iteration space, necessary for
+ // GPU fused schedule and no-loop optimization
+ if (canEmitGPUFusedDistSchedule(CGM, S, S.getDirectiveKind()) ||
+ NoLoopKernel) {
JumpDest LoopExit =
getJumpDestInCurrentScope(createBasicBlock("omp.loop.exit"));
CodeGenLoop(*this, S, LoopExit);
@@ -6421,10 +6502,13 @@ void CodeGenFunction::EmitOMPDistributeLoop(const OMPLoopDirective &S,
}
}
if (isOpenMPSimdDirective(S.getDirectiveKind())) {
- EmitOMPSimdFinal(S, [IL, &S](CodeGenFunction &CGF) {
- return CGF.Builder.CreateIsNotNull(
- CGF.EmitLoadOfScalar(IL, S.getBeginLoc()));
- });
+ EmitOMPSimdFinal(
+ S, [IL, &S, NoLoopKernel](CodeGenFunction &CGF) -> llvm::Value * {
+ if (NoLoopKernel)
+ return nullptr;
+ return CGF.Builder.CreateIsNotNull(
+ CGF.EmitLoadOfScalar(IL, S.getBeginLoc()));
+ });
}
if (isOpenMPSimdDirective(S.getDirectiveKind()) &&
!isOpenMPParallelDirective(S.getDirectiveKind()) &&
diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp
index fe39235fcd4f63..fa4404c66ca709 100644
--- a/clang/lib/CodeGen/CodeGenFunction.cpp
+++ b/clang/lib/CodeGen/CodeGenFunction.cpp
@@ -104,16 +104,18 @@ CodeGenFunction::~CodeGenFunction() {
assert(DeferredDeactivationCleanupStack.empty() &&
"missed to deactivate a cleanup");
- if (getLangOpts().OpenMP && CurFn)
+ if (getLangOpts().OpenMP && CurFn) {
CGM.getOpenMPRuntime().functionFinished(*this);
- // If we have an OpenMPIRBuilder we want to finalize functions (incl.
- // outlining etc) at some point. Doing it once the function codegen is done
- // seems to be a reasonable spot. We do it here, as opposed to the deletion
- // time of the CodeGenModule, because we have to ensure the IR has not yet
- // been "emitted" to the outside, thus, modifications are still sensible.
- if (CGM.getLangOpts().OpenMPIRBuilder && CurFn)
- CGM.getOpenMPRuntime().getOMPBuilder().finalize(CurFn);
+ // Finalizing (incl. outlining etc) once the function codegen is done, as
+ // opposed to the deletion time of the CodeGenModule, ensures the IR has
+ // not yet been "emitted" to the outside, thus, modifications are still
+ // sensible.
+ llvm::OpenMPIRBuilder &OMPBuilder = CGM.getOpenMPRuntime().getOMPBuilder();
+ if (CGM.getLangOpts().OpenMPIRBuilder ||
+ OMPBuilder.hasPendingOutlines(CurFn))
+ OMPBuilder.finalize(CurFn);
+ }
}
// Map the LangOption for exception behavior into
diff --git a/clang/lib/CodeGen/CodeGenFunction.h b/clang/lib/CodeGen/CodeGenFunction.h
index 3c8188c4cefdd7..26917a6dc0a8c4 100644
--- a/clang/lib/CodeGen/CodeGenFunction.h
+++ b/clang/lib/CodeGen/CodeGenFunction.h
@@ -4131,7 +4131,8 @@ class CodeGenFunction : public CodeGenTypeCache {
JumpDest getOMPCancelDestination(OpenMPDirectiveKind Kind);
/// Emit initial code for loop counters of loop-based directives.
void EmitOMPPrivateLoopCounters(const OMPLoopDirective &S,
- OMPPrivateScope &LoopScope);
+ OMPPrivateScope &LoopScope,
+ bool OnlyUnresolved = false);
/// Helper for the OpenMP loop directives.
void EmitOMPLoopBody(const OMPLoopDirective &D, JumpDest LoopExit);
diff --git a/clang/test/OpenMP/target_no_loop.c b/clang/test/OpenMP/target_no_loop.c
new file mode 100644
index 00000000000000..1caffb04a8c767
--- /dev/null
+++ b/clang/test/OpenMP/target_no_loop.c
@@ -0,0 +1,81 @@
+// REQUIRES: amdgpu-registered-target
+
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple x86_64-unknown-linux-gnu \
+// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-host.bc
+
+// 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: -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 \
+// RUN: -emit-llvm %s -o - | FileCheck %s \
+// RUN: --check-prefix=SPMD --implicit-check-not=__kmpc_distribute_for_static_loop_4u
+
+// 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: -emit-llvm %s -o - | FileCheck %s \
+// RUN: --check-prefix=SPMD --implicit-check-not=__kmpc_distribute_for_static_loop_4u
+
+// 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-threads-oversubscription \
+// RUN: -emit-llvm %s -o - | FileCheck %s \
+// RUN: --check-prefix=SPMD --implicit-check-not=__kmpc_distribute_for_static_loop_4u
+
+// expected-no-diagnostics
+
+void no_loop(int *array) {
+#pragma omp target teams distribute parallel for
+ for (int i = 0; i < 1024; ++i)
+ array[i] = i + 1;
+}
+
+void no_loop_simd(int *array) {
+#pragma omp target teams distribute parallel for simd
+ for (int i = 0; i < 1024; ++i)
+ array[i] = i + 1;
+}
+
+void no_loop_nowait(int *array) {
+#pragma omp target teams distribute parallel for nowait
+ for (int i = 0; i < 1024; ++i)
+ array[i] = i + 1;
+}
+
+// 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-LABEL: @__kmpc_parallel_60({{.*}}no_loop_l{{[0-9]+}}{{.*}})
+// NOLOOP: omp.loop.exit:
+// NOLOOP-NEXT: ret void
+// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}no_loop_l{{[0-9]+}}{{.*}}, i32 0, i32 0, i8 1)
+// NOLOOP: omp_loop.after:
+// NOLOOP-NEXT: ret void
+
+// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}no_loop_simd{{.*}})
+// NOLOOP: omp.loop.exit:
+// NOLOOP-NEXT: store i32 1024, ptr %i
+// NOLOOP-NEXT: ret void
+// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}no_loop_simd{{.*}}, i32 0, i32 0, i8 1)
+// NOLOOP: omp_loop.after:
+// NOLOOP-NEXT: ret void
+
+// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}no_loop_nowait{{.*}})
+// NOLOOP: omp.loop.exit:
+// NOLOOP-NEXT: ret void
+// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}no_loop_nowait{{.*}}, 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
+
+// SPMD-COUNT-3: _kernel_environment {{.*}} i8 0, i8 1, i8 2
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
index a13832d69b6a09..ccf01e504269fe 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
@@ -2731,6 +2731,14 @@ class OpenMPIRBuilder {
OutlineInfos.emplace_back(std::move(OI));
}
+ /// Return true if \p Fn has a region registered for outlining that has not
+ /// been processed yet.
+ bool hasPendingOutlines(const Function *Fn) const {
+ return any_of(OutlineInfos, [Fn](const std::unique_ptr<OutlineInfo> &OI) {
+ return OI->getFunction() == Fn;
+ });
+ }
+
/// An ordered map of auto-generated variables to their unique names.
/// It stores variables with the following names: 1) ".gomp_critical_user_" +
/// <critical_section_name> + ".var" for "omp critical" directives; 2)
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 8b196c461f195b..4fee93b9d90a4e 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -1063,7 +1063,7 @@ void OpenMPIRBuilder::finalize(Function *Fn) {
"OMPIRBuilder finalization \n";
};
- if (!OffloadInfoManager.empty())
+ if (!Fn && !OffloadInfoManager.empty())
createOffloadEntriesAndInfoMetadata(ErrorReportFn);
// Rewrite uses of globals to their replacement declare target globals if
diff --git a/offload/test/offloading/target-no-loop.c b/offload/test/offloading/target-no-loop.c
new file mode 100644
index 00000000000000..59558c67152e25
--- /dev/null
+++ b/offload/test/offloading/target-no-loop.c
@@ -0,0 +1,91 @@
+// clang-format off
+// C counterpart of fortran/target-no-loop.f90.
+
+// RUN: %libomptarget-compile-generic -O3 -fopenmp-assume-threads-oversubscription -fopenmp-assume-teams-oversubscription
+// RUN: env LIBOMPTARGET_INFO=16 OMP_NUM_TEAMS=16 OMP_TEAMS_THREAD_LIMIT=16 %libomptarget-run-generic 2>&1 | %fcheck-generic
+// REQUIRES: gpu
+// XFAIL: intelgpu
+
+#include <stdio.h>
+
+static int check_errors(int *array) {
+ int errors = 0;
+ for (int i = 0; i < 1024; ++i)
+ if (array[i] != i + 1)
+ ++errors;
+ return errors;
+}
+
+int main(void) {
+ int array[1024];
+ int errors = 0;
+ int red;
+
+ for (int i = 0; i < 1024; ++i)
+ array[i] = 1;
+
+ // No-loop kernel
+#pragma omp target teams distribute parallel for
+ for (int i = 0; i < 1024; ++i)
+ array[i] = i + 1;
+ errors += check_errors(array);
+
+ // SPMD kernel (num_teams clause blocks promotion to no-loop)
+ for (int i = 0; i < 1024; ++i)
+ array[i] = 1;
+#pragma omp target teams distribute parallel for num_teams(3)
+ for (int i = 0; i < 1024; ++i)
+ array[i] = i + 1;
+ errors += check_errors(array);
+
+ // No-loop kernel
+ for (int i = 0; i < 1024; ++i)
+ array[i] = 1;
+#pragma omp target teams distribute parallel for num_threads(64)
+ for (int i = 0; i < 1024; ++i)
+ array[i] = i + 1;
+ errors += check_errors(array);
+
+ // SPMD kernel
+ for (int i = 0; i < 1024; ++i)
+ array[i] = 1;
+#pragma omp target parallel for
+ for (int i = 0; i < 1024; ++i)
+ array[i] = i + 1;
+ errors += check_errors(array);
+
+ // Generic kernel
+ for (int i = 0; i < 1024; ++i)
+ array[i] = 1;
+#pragma omp target teams distribute
+ for (int i = 0; i < 1024; ++i)
+ array[i] = i + 1;
+ errors += check_errors(array);
+
+ // SPMD kernel (reduction clause blocks promotion to no-loop)
+ for (int i = 0; i < 1024; ++i)
+ array[i] = 1;
+ red = 0;
+#pragma omp target teams distribute parallel for reduction(+ : red)
+ for (int i = 0; i < 1024; ++i)
+ red += array[i];
+ if (red != 1024)
+ ++errors;
+
+ printf("number of errors: %d\n", errors);
+ return 0;
+}
+
+// CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}} SPMD-No-Loop mode
+// CHECK: info: #Args: 2 Teams x Thrds: 64x 16
+// CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}} SPMD mode
+// CHECK: info: #Args: 2 Teams x Thrds: 3x 16 {{.*}}
+// CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}} SPMD-No-Loop mode
+// CHECK: info: #Args: 2 Teams x Thrds: 64x 16 {{.*}}
+// CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}} SPMD mode
+// CHECK: info: #Args: 2 Teams x Thrds: 1x 16
+// CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}} Generic-SPMD mode
+// CHECK: info: #Args: 2 Teams x Thrds: 16x 16 {{.*}}
+// CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}} SPMD mode
+// CHECK: info: #Args: 3 Teams x Thrds: 16x 16 {{.*}}
+// CHECK: number of errors: 0
>From 7ff0f54fbabb69638ee3ab417b55a483b2ef5264 Mon Sep 17 00:00:00 2001
From: Nicole Aschenbrenner <nicole.aschenbrenner at amd.com>
Date: Fri, 25 Sep 2026 04:27:56 -0500
Subject: [PATCH 2/2] Remove unused OnlyUnresolved parameter
---
clang/lib/CodeGen/CGStmtOpenMP.cpp | 18 ++----------------
clang/lib/CodeGen/CodeGenFunction.h | 3 +--
2 files changed, 3 insertions(+), 18 deletions(-)
diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp
index 2836911518398e..f257c38d84b900 100644
--- a/clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -2668,26 +2668,12 @@ static void emitAlignedClause(CodeGenFunction &CGF,
}
void CodeGenFunction::EmitOMPPrivateLoopCounters(
- const OMPLoopDirective &S, CodeGenFunction::OMPPrivateScope &LoopScope,
- bool OnlyUnresolved) {
+ const OMPLoopDirective &S, CodeGenFunction::OMPPrivateScope &LoopScope) {
if (!HaveInsertPoint())
return;
auto I = S.private_counters().begin();
for (const Expr *E : S.counters()) {
- const auto *DRE = cast<DeclRefExpr>(E);
- const auto *VD = cast<VarDecl>(DRE->getDecl());
- // Skip counters that already resolve, mirroring EmitDeclRefLValue's
- // handling for these cases.
- if (OnlyUnresolved) {
- const VarDecl *Canonical = VD->getCanonicalDecl();
- if (!DRE->refersToEnclosingVariableOrCapture() || !CapturedStmtInfo ||
- LocalDeclMap.count(Canonical) ||
- CapturedStmtInfo->lookup(Canonical)) {
- ++I;
- continue;
- }
- }
-
+ const auto *VD = cast<VarDecl>(cast<DeclRefExpr>(E)->getDecl());
const auto *PrivateVD = cast<VarDecl>(cast<DeclRefExpr>(*I)->getDecl());
// Emit var without initialization.
AutoVarEmission VarEmission = EmitAutoVarAlloca(*PrivateVD);
diff --git a/clang/lib/CodeGen/CodeGenFunction.h b/clang/lib/CodeGen/CodeGenFunction.h
index 26917a6dc0a8c4..3c8188c4cefdd7 100644
--- a/clang/lib/CodeGen/CodeGenFunction.h
+++ b/clang/lib/CodeGen/CodeGenFunction.h
@@ -4131,8 +4131,7 @@ class CodeGenFunction : public CodeGenTypeCache {
JumpDest getOMPCancelDestination(OpenMPDirectiveKind Kind);
/// Emit initial code for loop counters of loop-based directives.
void EmitOMPPrivateLoopCounters(const OMPLoopDirective &S,
- OMPPrivateScope &LoopScope,
- bool OnlyUnresolved = false);
+ OMPPrivateScope &LoopScope);
/// Helper for the OpenMP loop directives.
void EmitOMPLoopBody(const OMPLoopDirective &D, JumpDest LoopExit);
More information about the llvm-commits
mailing list