[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 20 03:25:22 PDT 2026
https://github.com/nicebert updated https://github.com/llvm/llvm-project/pull/224045
>From 7636b888556468c2135c3a115e256aae1a6c980b 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 1/2] [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 | 56 +++++++++++++-
clang/test/OpenMP/target_no_loop.cpp | 96 +++++++++++++++++++-----
offload/test/offloading/target-no-loop.c | 48 ++++++++++++
4 files changed, 176 insertions(+), 26 deletions(-)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
index ecb77ac97bb675..97c43734717642 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
@@ -721,8 +721,6 @@ checkAndGetNoLoopDirective(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 nullptr;
diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp
index 0e1d839ff04e33..b59515cbde56ed 100644
--- a/clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -74,6 +74,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.
@@ -3972,6 +3995,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);
@@ -3995,17 +4022,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;
}
@@ -6592,7 +6635,9 @@ void CodeGenFunction::EmitOMPDistributeLoop(const OMPLoopDirective &S,
EmitOMPReductionClauseInit(S, LoopScope);
const bool NoLoopKernel = CGM.getOpenMPRuntime().canPromoteToNoLoop(S);
- HasLastprivateClause = EmitOMPLastprivateClauseInit(S, LoopScope);
+ if (!NoLoopKernel)
+ HasLastprivateClause = EmitOMPLastprivateClauseInit(S, LoopScope);
+
EmitOMPPrivateLoopCounters(S, LoopScope);
(void)LoopScope.Privatize();
if (isOpenMPTargetExecutionDirective(S.getDirectiveKind()))
@@ -6728,7 +6773,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 1eb24fe6d2542e..cb8953bc3588c6 100644
--- a/clang/test/OpenMP/target_no_loop.cpp
+++ b/clang/test/OpenMP/target_no_loop.cpp
@@ -79,6 +79,38 @@ void promotable_no_loop(int *array, int n) {
#pragma omp target teams distribute parallel for if(1)
for (int i = 0; i < 1024; ++i)
array[i] = i + 1;
+
+ {
+ 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, int n) {
@@ -105,22 +137,6 @@ void non_promotable(int *array, int n) {
#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;
- }
-
#pragma omp target teams distribute parallel for if(parallel: n > 0)
for (int i = 0; i < 1024; ++i)
array[i] = i + 1;
@@ -145,13 +161,13 @@ void promotable_after_reduction(int *array) {
array[i] = i + offset;
}
-// NOLOOP-COUNT-7: promotable_no_loop{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 6
-// NOLOOP-COUNT-9: non_promotable{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 2
+// NOLOOP-COUNT-11: 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: promotable_after_reduction{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 2
// NOLOOP: promotable_after_reduction{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 6
-// SPMD-COUNT-7: promotable_no_loop{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 2
-// SPMD-COUNT-9: non_promotable{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 2
+// SPMD-COUNT-11: 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-2: promotable_after_reduction{{.*}}_kernel_environment {{.*}} i8 0, i8 1, i8 2
// reduction before a promoted kernel
@@ -227,3 +243,43 @@ void promotable_after_reduction(int *array) {
// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l74_{{.*}}, 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{{.*}}_l85_{{.*}})
+// 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({{.*}}_l85_{{.*}}, 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{{.*}}_l92_{{.*}})
+// NOLOOP: omp.loop.exit:
+// NOLOOP-NEXT: @__kmpc_free_shared(ptr %last{{.*}})
+// NOLOOP-NEXT: ret void
+// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l92_{{.*}}, 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{{.*}}_l101_{{.*}})
+// NOLOOP: omp.loop.exit:
+// NOLOOP-NEXT: @__kmpc_free_shared(ptr %last{{.*}})
+// NOLOOP-NEXT: ret void
+// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l101_{{.*}}, 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{{.*}}_l110_{{.*}})
+// NOLOOP: omp.loop.exit:
+// NOLOOP-NEXT: store i32 1024, ptr %i
+// NOLOOP-NEXT: ret void
+// NOLOOP: @__kmpc_distribute_for_static_loop_4u({{.*}}_l110_{{.*}}, 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 8407cffc13020f..a3149025044c31 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
>From b9c0e9b74b84d05f619124f3bf12e15e4418812c Mon Sep 17 00:00:00 2001
From: Nicole Aschenbrenner <nicole.aschenbrenner at amd.com>
Date: Wed, 14 Oct 2026 09:09:06 -0500
Subject: [PATCH 2/2] Add braces around loop body with nested if
---
clang/lib/CodeGen/CGStmtOpenMP.cpp | 5 +++--
1 file changed, 3 insertions(+), 2 deletions(-)
diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp
index b59515cbde56ed..4cd348bbb5f6ad 100644
--- a/clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -78,10 +78,11 @@ 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())
+ for (const Expr *C : S.counters()) {
if (cast<VarDecl>(cast<DeclRefExpr>(C)->getDecl())->getCanonicalDecl() ==
Canonical)
return true;
+ }
return false;
}
@@ -6637,7 +6638,7 @@ void CodeGenFunction::EmitOMPDistributeLoop(const OMPLoopDirective &S,
const bool NoLoopKernel = CGM.getOpenMPRuntime().canPromoteToNoLoop(S);
if (!NoLoopKernel)
HasLastprivateClause = EmitOMPLastprivateClauseInit(S, LoopScope);
-
+
EmitOMPPrivateLoopCounters(S, LoopScope);
(void)LoopScope.Privatize();
if (isOpenMPTargetExecutionDirective(S.getDirectiveKind()))
More information about the llvm-branch-commits
mailing list