[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
Thu Sep 17 04:17:44 PDT 2026
https://github.com/nicebert updated https://github.com/llvm/llvm-project/pull/224045
>From a02b869b1a4a29e122ae76271ce69594711c6d88 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 | 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 264d78b0f5245..a22ccfcf3ddc2 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
@@ -1158,9 +1158,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 2836911518398..4c59c726c66ff 100644
--- a/clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -66,6 +66,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);
@@ -3772,6 +3795,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);
@@ -3795,17 +3822,33 @@ emitInnerParallelForWhenCombined(CodeGenFunction &CGF,
CGF.AllocaInsertPt->getParent(), 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;
}
@@ -6364,8 +6407,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()))
@@ -6501,7 +6547,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 59558c67152e2..ae6c31c966b19 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: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}} SPMD mode
// CHECK: info: #Args: 3 Teams x Thrds: 16x 16 {{.*}}
+// CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}} SPMD-No-Loop mode
+// CHECK: info: #Args: 3 Teams x Thrds: 64x 16 {{.*}}
+// CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}} SPMD-No-Loop mode
+// CHECK: info: #Args: 3 Teams x Thrds: 64x 16 {{.*}}
+// CHECK: PluginInterface 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