[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
Wed Oct 7 05:25:15 PDT 2026
https://github.com/nicebert updated https://github.com/llvm/llvm-project/pull/224045
>From 687b07ac133a344e5716abd8be35fbc70f270f55 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] [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 | 2 -
clang/lib/CodeGen/CGStmtOpenMP.cpp | 57 +++++++++++++-
clang/test/OpenMP/target_no_loop.cpp | 96 +++++++++++++++++++-----
offload/test/offloading/target-no-loop.c | 48 ++++++++++++
4 files changed, 177 insertions(+), 26 deletions(-)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
index af342eaec2da4..a860aace68b39 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
@@ -720,8 +720,6 @@ static bool isNoLoopEligible(ASTContext &Ctx, const OMPExecutableDirective &D) {
// Do not filter out 'ordered' since Sema rejects it on these directives.
if (LD.getLoopsNumber() != 1 || LD.hasClausesOfKind<OMPNumTeamsClause>() ||
LD.hasClausesOfKind<OMPReductionClause>() ||
- LD.hasClausesOfKind<OMPLastprivateClause>() ||
- LD.hasClausesOfKind<OMPLinearClause>() ||
LD.getSingleClause<OMPScheduleClause>() ||
LD.getSingleClause<OMPDistScheduleClause>())
return false;
diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp
index 9a5ab88a91f6b..0f0802de3d19a 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;
+}
+
namespace {
/// Lexical scope for OpenMP executable constructs, that handles correct codegen
/// for captured expressions.
@@ -3973,6 +3996,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);
@@ -3996,17 +4023,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;
}
@@ -6594,8 +6637,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 = CGM.getOpenMPRuntime().canPromoteToNoLoop();
- HasLastprivateClause = EmitOMPLastprivateClauseInit(S, LoopScope);
+ if (!NoLoopKernel)
+ HasLastprivateClause = EmitOMPLastprivateClauseInit(S, LoopScope);
+
EmitOMPPrivateLoopCounters(S, LoopScope);
(void)LoopScope.Privatize();
if (isOpenMPTargetExecutionDirective(S.getDirectiveKind()))
@@ -6731,7 +6777,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.cpp b/clang/test/OpenMP/target_no_loop.cpp
index 1859b26ae3fb8..ca21a8793ab8b 100644
--- a/clang/test/OpenMP/target_no_loop.cpp
+++ b/clang/test/OpenMP/target_no_loop.cpp
@@ -65,6 +65,38 @@ void promotable_no_loop(int *array, int n) {
for (int i = 0; i < 1024; ++i)
set(i);
}
+
+ {
+ int i;
+#pragma omp target teams distribute parallel for lastprivate(i)
+ for (i = 0; i < 1024; ++i)
+ array[i] = i + 1;
+ }
+
+ {
+ 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;
+ }
+ }
+
+ {
+ 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;
+ }
+ }
+
+ {
+ int i;
+#pragma omp target teams distribute parallel for simd linear(i)
+ for (i = 0; i < 1024; ++i)
+ array[i] = i + 1;
+ }
}
void non_promotable(int *array) {
@@ -90,29 +122,13 @@ void non_promotable(int *array) {
array[i] = i + 1;
#pragma omp cancel for
}
-
- {
- 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;
- }
- }
-
- {
- int i;
-#pragma omp target teams distribute parallel for simd linear(i)
- for (i = 0; i < 1024; ++i)
- array[i] = i + 1;
- }
}
-// NOLOOP-COUNT-6: promotable_no_loop{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 6
-// NOLOOP-COUNT-7: non_promotable{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 2
+// NOLOOP-COUNT-10: promotable_no_loop{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 6
+// NOLOOP-COUNT-5: non_promotable{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 2
-// SPMD-COUNT-6: promotable_no_loop{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 2
-// SPMD-COUNT-7: non_promotable{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 2
+// SPMD-COUNT-10: promotable_no_loop{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 2
+// SPMD-COUNT-5: non_promotable{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 2
// no clause
// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l37_{{.*}})
@@ -171,3 +187,43 @@ void non_promotable(int *array) {
// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l64_{{.*}}, i32 0, i32 0, i8 1)
// NOLOOP: omp_loop.after:
// NOLOOP-NEXT: ret void
+
+// lastprivate loop counter
+// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l71_{{.*}})
+// NOLOOP: omp.loop.exit:
+// NOLOOP-NEXT: store i32 1024, ptr %i
+// NOLOOP-NEXT: @__kmpc_free_shared(ptr %i{{.*}})
+// NOLOOP-NEXT: ret void
+// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l71_{{.*}}, i32 0, i32 0, i8 1)
+// NOLOOP: @__kmpc_barrier
+// NOLOOP: omp_loop.after:
+// NOLOOP-NEXT: ret void
+
+// lastprivate scalar
+// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l78_{{.*}})
+// NOLOOP: omp.loop.exit:
+// NOLOOP-NEXT: @__kmpc_free_shared(ptr %last{{.*}})
+// NOLOOP-NEXT: ret void
+// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l78_{{.*}}, i32 0, i32 0, i8 1)
+// NOLOOP: @__kmpc_barrier
+// NOLOOP: store {{.*}}, ptr %last.
+// NOLOOP-NEXT: %.omp.lastprivate.done
+
+// lastprivate scalar, nowait
+// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l87_{{.*}})
+// NOLOOP: omp.loop.exit:
+// NOLOOP-NEXT: @__kmpc_free_shared(ptr %last{{.*}})
+// NOLOOP-NEXT: ret void
+// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l87_{{.*}}, i32 0, i32 0, i8 1)
+// NOLOOP: @__kmpc_barrier
+// NOLOOP: store {{.*}}, ptr %last.
+// NOLOOP-NEXT: %.omp.lastprivate.done
+
+// simd linear loop counter
+// NOLOOP-LABEL: @__kmpc_parallel_60({{.*}}promotable_no_loop{{.*}}_l96_{{.*}})
+// NOLOOP: omp.loop.exit:
+// NOLOOP-NEXT: store i32 1024, ptr %i
+// NOLOOP-NEXT: ret void
+// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l96_{{.*}}, i32 0, i32 0, i8 1)
+// NOLOOP: omp_loop.after:
+// NOLOOP-NEXT: ret void
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