[llvm-branch-commits] [clang] [llvm] [clang][OpenMP] Emit lastprivate final copies in no-loop kernels (PR #224045)
Nicole Aschenbrenner via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Tue Oct 6 05:52:22 PDT 2026
https://github.com/nicebert updated https://github.com/llvm/llvm-project/pull/224045
>From c709890f9f3bb2b42beffcc14d9fd30c54075ed4 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/3] [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 ae55295cadc14..9a93f5c5c2e82 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 e7322bdd0deca..12e4f718d64ed 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
@@ -763,9 +763,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;
@@ -1158,6 +1163,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 50efbfd1577a5..d40d15767552e 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.h
+++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.h
@@ -146,6 +146,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 0913f76018a58..1f57c8292e5b2 100644
--- a/clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -75,6 +75,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.
@@ -2574,6 +2583,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");
@@ -2649,11 +2676,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();
@@ -2664,14 +2687,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.
@@ -2861,12 +2881,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);
@@ -3915,6 +3949,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,
@@ -3934,6 +3984,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->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);
@@ -4012,19 +4111,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();
@@ -4119,6 +4206,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.
@@ -6469,19 +6557,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();
@@ -6540,6 +6616,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();
@@ -6562,12 +6640,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);
@@ -6674,10 +6755,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 62dade6cf6f49..78e37b7b2b853 100644
--- a/clang/lib/CodeGen/CodeGenFunction.cpp
+++ b/clang/lib/CodeGen/CodeGenFunction.cpp
@@ -88,16 +88,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 490267aaefd86..caefc275771ff 100644
--- a/clang/lib/CodeGen/CodeGenFunction.h
+++ b/clang/lib/CodeGen/CodeGenFunction.h
@@ -4185,7 +4185,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 0000000000000..1caffb04a8c76
--- /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 5acbb69248309..cfc84226b34a7 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
@@ -2734,6 +2734,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 b6a11bc3c805e..4f98f59d71e31 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 0000000000000..8407cffc13020
--- /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: omptarget device {{[0-9]+}} info: Launching kernel {{.*}} SPMD-No-Loop mode
+// CHECK: info: #Args: 2 Teams x Thrds: 64x 16
+// CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}} SPMD mode
+// CHECK: info: #Args: 2 Teams x Thrds: 3x 16 {{.*}}
+// CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}} SPMD-No-Loop mode
+// CHECK: info: #Args: 2 Teams x Thrds: 64x 16 {{.*}}
+// CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}} SPMD mode
+// CHECK: info: #Args: 2 Teams x Thrds: 1x 16
+// CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}} Generic-SPMD mode
+// CHECK: info: #Args: 2 Teams x Thrds: 16x 16 {{.*}}
+// CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}} SPMD mode
+// CHECK: info: #Args: 3 Teams x Thrds: 16x 16 {{.*}}
+// CHECK: number of errors: 0
>From a38f7477e942ccd25165b4aef574859fdea3a716 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/3] 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 1f57c8292e5b2..251241b38f918 100644
--- a/clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -2881,26 +2881,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 caefc275771ff..490267aaefd86 100644
--- a/clang/lib/CodeGen/CodeGenFunction.h
+++ b/clang/lib/CodeGen/CodeGenFunction.h
@@ -4185,8 +4185,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);
>From 456cf124632b959d522d484c5ebcda19a04d58c4 Mon Sep 17 00:00:00 2001
From: Nicole Aschenbrenner <nicole.aschenbrenner at amd.com>
Date: Wed, 16 Sep 2026 05:35:29 -0500
Subject: [PATCH 3/3] [clang][OpenMP] Emit lastprivate final copies in no-loop
kernels
Promotion rejected lastprivate because the no-loop branch leaves the
worksharing path before it privatizes or copies out, so the clause would
have been silently dropped.
Privatize the non-counter variables in the parallel region and copy them
out under the last iteration that applyWorkshareLoop publishes, forcing
the exit barrier the copy reads through. Loop counters stay at the
distribute level and reach EmitOMPSimdFinal.
---
clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp | 4 +-
clang/lib/CodeGen/CGStmtOpenMP.cpp | 57 ++++++++++++++++++++++--
clang/test/OpenMP/target_no_loop.c | 55 ++++++++++++++++++++++-
offload/test/offloading/target-no-loop.c | 48 ++++++++++++++++++++
4 files changed, 156 insertions(+), 8 deletions(-)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
index 12e4f718d64ed..fa6fcdab47b18 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
@@ -1171,9 +1171,7 @@ bool CGOpenMPRuntimeGPU::canPromoteToNoLoop(
DKind == OMPD_target_teams_distribute_parallel_for_simd) &&
LangOpts.OpenMPTeamSubscription && LangOpts.OpenMPThreadSubscription &&
!D.hasClausesOfKind<OMPNumTeamsClause>() &&
- !D.hasClausesOfKind<OMPReductionClause>() &&
- !D.hasClausesOfKind<OMPLastprivateClause>() &&
- !D.hasClausesOfKind<OMPLinearClause>();
+ !D.hasClausesOfKind<OMPReductionClause>();
}
std::pair<llvm::Value *, llvm::Value *>
diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp
index 251241b38f918..1a4e507e06390 100644
--- a/clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -75,6 +75,29 @@ static bool canEmitGPUFusedDistSchedule(const CodeGenModule &CGM,
!S.getSingleClause<OMPOrderedClause>();
}
+static bool isLoopCounter(const OMPLoopDirective &S, const Expr *E) {
+ const VarDecl *Canonical =
+ cast<VarDecl>(cast<DeclRefExpr>(E->IgnoreParenImpCasts())->getDecl())
+ ->getCanonicalDecl();
+ for (const Expr *C : S.counters())
+ if (cast<VarDecl>(cast<DeclRefExpr>(C)->getDecl())->getCanonicalDecl() ==
+ Canonical)
+ return true;
+ return false;
+}
+
+static bool
+hasLoopLastprivateVariable(const OMPLoopDirective &S,
+ llvm::function_ref<bool(const Expr *)> Pred) {
+ for (const auto *C : S.getClausesOfKind<OMPLastprivateClause>()) {
+ for (const Expr *E : C->varlist()) {
+ if (Pred(E))
+ return true;
+ }
+ }
+ return false;
+}
+
static bool canEmitGPUNoLoopKernel(CodeGenModule &CGM,
const OMPLoopDirective &S) {
const auto *D = dyn_cast<OMPTargetTeamsDistributeParallelForDirective>(&S);
@@ -3982,6 +4005,10 @@ emitInnerParallelForWhenCombined(CodeGenFunction &CGF,
CGF.EmitOMPPrivateLoopCounters(S, PrivateScope);
(void)PrivateScope.Privatize();
+ CodeGenFunction::OMPPrivateScope LastprivateScope(CGF);
+ (void)CGF.EmitOMPLastprivateClauseInit(S, LastprivateScope);
+ (void)LastprivateScope.Privatize();
+
if (isOpenMPTargetExecutionDirective(EKind))
CGM.getOpenMPRuntime().adjustTargetSpecificDataForLambdas(CGF, S);
@@ -4005,17 +4032,33 @@ emitInnerParallelForWhenCombined(CodeGenFunction &CGF,
CGF.AllocaInsertPt->getIterator());
llvm::OpenMPIRBuilder &OMPBuilder =
CGM.getOpenMPRuntime().getOMPBuilder();
+ bool NeedsLastprivateFinalCopy = hasLoopLastprivateVariable(
+ S, [&S](const Expr *E) { return !isLoopCounter(S, E); });
cantFail(OMPBuilder.applyWorkshareLoop(
CGF.Builder.getCurrentDebugLocation(), CLI, AllocaIP,
- /*NeedsBarrier=*/!S.getSingleClause<OMPNowaitClause>(),
+ /*NeedsBarrier=*/!S.getSingleClause<OMPNowaitClause>() ||
+ NeedsLastprivateFinalCopy,
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));
+ /*DistScheduleChunkSize=*/nullptr,
+ /*NeedsLastIter=*/NeedsLastprivateFinalCopy));
+
+ // Emit final copy for lastprivate variables, excluding counters deferred
+ // to the distribute level.
+ if (NeedsLastprivateFinalCopy) {
+ llvm::Value *LastIter = CLI->getLastIter();
+ assert(LastIter && "workshare loop did not publish last iteration");
+ CGF.EmitOMPLastprivateClauseFinal(
+ S, /*NoFinals=*/true,
+ CGF.Builder.CreateIsNotNull(CGF.Builder.CreateAlignedLoad(
+ CGF.Int32Ty, LastIter, CharUnits::fromQuantity(4),
+ ".omp.is_last")));
+ }
return;
}
@@ -6603,8 +6646,11 @@ void CodeGenFunction::EmitOMPDistributeLoop(const OMPLoopDirective &S,
!isOpenMPTeamsDirective(S.getDirectiveKind()))
EmitOMPReductionClauseInit(S, LoopScope);
+ // Defer lastprivate init to the parallel region for no-loop
const bool NoLoopKernel = canEmitGPUNoLoopKernel(CGM, S);
- HasLastprivateClause = EmitOMPLastprivateClauseInit(S, LoopScope);
+ if (!NoLoopKernel)
+ HasLastprivateClause = EmitOMPLastprivateClauseInit(S, LoopScope);
+
EmitOMPPrivateLoopCounters(S, LoopScope);
(void)LoopScope.Privatize();
if (isOpenMPTargetExecutionDirective(S.getDirectiveKind()))
@@ -6740,7 +6786,10 @@ void CodeGenFunction::EmitOMPDistributeLoop(const OMPLoopDirective &S,
CodeGenLoop);
}
}
- if (isOpenMPSimdDirective(S.getDirectiveKind())) {
+ if (isOpenMPSimdDirective(S.getDirectiveKind()) ||
+ (NoLoopKernel && hasLoopLastprivateVariable(S, [&S](const Expr *E) {
+ return isLoopCounter(S, E);
+ }))) {
EmitOMPSimdFinal(
S, [IL, &S, NoLoopKernel](CodeGenFunction &CGF) -> llvm::Value * {
if (NoLoopKernel)
diff --git a/clang/test/OpenMP/target_no_loop.c b/clang/test/OpenMP/target_no_loop.c
index 1caffb04a8c76..41b7ab96e8867 100644
--- a/clang/test/OpenMP/target_no_loop.c
+++ b/clang/test/OpenMP/target_no_loop.c
@@ -45,12 +45,37 @@ void no_loop_simd(int *array) {
array[i] = i + 1;
}
+void no_loop_lastprivate_counter(int *array) {
+ int i;
+#pragma omp target teams distribute parallel for lastprivate(i)
+ for (i = 0; i < 1024; ++i)
+ array[i] = i + 1;
+}
+
+void no_loop_lastprivate_scalar(int *array) {
+ int last = 0;
+#pragma omp target teams distribute parallel for lastprivate(last)
+ for (int i = 0; i < 1024; ++i) {
+ array[i] = i + 1;
+ last = i;
+ }
+}
+
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;
}
+void no_loop_lastprivate_scalar_nowait(int *array) {
+ int last = 0;
+#pragma omp target teams distribute parallel for lastprivate(last) nowait
+ for (int i = 0; i < 1024; ++i) {
+ array[i] = i + 1;
+ last = i;
+ }
+}
+
// 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
@@ -69,6 +94,25 @@ void no_loop_nowait(int *array) {
// NOLOOP: omp_loop.after:
// NOLOOP-NEXT: ret void
+// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}lastprivate_counter{{.*}})
+// 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({{.*}}lastprivate_counter{{.*}}, i32 0, i32 0, i8 1)
+// NOLOOP: @__kmpc_barrier
+// NOLOOP: omp_loop.after:
+// NOLOOP-NEXT: ret void
+
+// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}lastprivate_scalar{{.*}})
+// NOLOOP: omp.loop.exit:
+// NOLOOP-NEXT: @__kmpc_free_shared(ptr %last{{.*}})
+// NOLOOP-NEXT: ret void
+// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}lastprivate_scalar{{.*}}, i32 0, i32 0, i8 1)
+// NOLOOP: @__kmpc_barrier
+// NOLOOP: store {{.*}}, ptr %last.
+// NOLOOP-NEXT: %.omp.lastprivate.done
+
// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}no_loop_nowait{{.*}})
// NOLOOP: omp.loop.exit:
// NOLOOP-NEXT: ret void
@@ -78,4 +122,13 @@ void no_loop_nowait(int *array) {
// NOLOOP: omp_loop.after:
// NOLOOP-NEXT: ret void
-// SPMD-COUNT-3: _kernel_environment {{.*}} i8 0, i8 1, i8 2
+// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}lastprivate_scalar_nowait{{.*}})
+// NOLOOP: omp.loop.exit:
+// NOLOOP-NEXT: @__kmpc_free_shared(ptr %last{{.*}})
+// NOLOOP-NEXT: ret void
+// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}lastprivate_scalar_nowait{{.*}}, i32 0, i32 0, i8 1)
+// NOLOOP: @__kmpc_barrier
+// NOLOOP: store {{.*}}, ptr %last.
+// NOLOOP-NEXT: %.omp.lastprivate.done
+
+// SPMD-COUNT-6: _kernel_environment {{.*}} i8 0, i8 1, i8 2
diff --git a/offload/test/offloading/target-no-loop.c b/offload/test/offloading/target-no-loop.c
index 8407cffc13020..a3149025044c3 100644
--- a/offload/test/offloading/target-no-loop.c
+++ b/offload/test/offloading/target-no-loop.c
@@ -72,6 +72,48 @@ int main(void) {
if (red != 1024)
++errors;
+ // No-loop kernel with a lastprivate variable (final value from thread executing last iteration)
+ for (int i = 0; i < 1024; ++i)
+ array[i] = 1;
+
+ int last = -1;
+#pragma omp target teams distribute parallel for lastprivate(last)
+ for (int i = 0; i < 1024; ++i) {
+ array[i] = i + 1;
+ last = i;
+ }
+ errors += check_errors(array);
+#ifdef ASSERT_LASTPRIVATE
+ if (last != 1023)
+ ++errors;
+#endif
+
+ // No-loop kernel with a lastprivate loop counter (final value from evaluation at distribution level)
+ for (int i = 0; i < 1024; ++i)
+ array[i] = 1;
+ int iv = -1;
+#pragma omp target teams distribute parallel for lastprivate(iv)
+ for (iv = 0; iv < 1024; ++iv)
+ array[iv] = iv + 1;
+ errors += check_errors(array);
+#ifdef ASSERT_LASTPRIVATE
+ if (iv != 1024)
+ ++errors;
+#endif
+
+ // No-loop simd kernel with a lastprivate loop counter
+ for (int i = 0; i < 1024; ++i)
+ array[i] = 1;
+ int simd_iv = -1;
+#pragma omp target teams distribute parallel for simd lastprivate(simd_iv)
+ for (simd_iv = 0; simd_iv < 1024; ++simd_iv)
+ array[simd_iv] = simd_iv + 1;
+ errors += check_errors(array);
+#ifdef ASSERT_LASTPRIVATE
+ if (simd_iv != 1024)
+ ++errors;
+#endif
+
printf("number of errors: %d\n", errors);
return 0;
}
@@ -88,4 +130,10 @@ int main(void) {
// CHECK: info: #Args: 2 Teams x Thrds: 16x 16 {{.*}}
// CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}} SPMD mode
// CHECK: info: #Args: 3 Teams x Thrds: 16x 16 {{.*}}
+// CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}} SPMD-No-Loop mode
+// CHECK: info: #Args: 3 Teams x Thrds: 64x 16 {{.*}}
+// CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}} SPMD-No-Loop mode
+// CHECK: info: #Args: 3 Teams x Thrds: 64x 16 {{.*}}
+// CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}} SPMD-No-Loop mode
+// CHECK: info: #Args: 3 Teams x Thrds: 64x 16 {{.*}}
// CHECK: number of errors: 0
More information about the llvm-branch-commits
mailing list