[clang] [llvm] [OpenMP][Clang] Add no-loop SPMD kernel promotion (PR #205325)
Nicole Aschenbrenner via cfe-commits
cfe-commits at lists.llvm.org
Wed Jul 29 07:10:21 PDT 2026
https://github.com/nicebert updated https://github.com/llvm/llvm-project/pull/205325
>From e1e11612988cf5672688f91d324e4975e6093e65 Mon Sep 17 00:00:00 2001
From: Nicole Aschenbrenner <nicole.aschenbrenner at amd.com>
Date: Fri, 19 Jun 2026 05:50:00 -0500
Subject: [PATCH 1/3] [OpenMP][Clang] Add no-loop SPMD kernel promotion
Promote target teams distribute parallel for kernels to no-loop SPMD
mode when teams and threads oversubscription are assumed and no
num_teams or reduction clause is present, mirroring Flang's MLIR
promotion. Adds a C offload test alongside the existing Fortran one.
---
clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp | 23 +++++-
offload/test/offloading/target-no-loop.c | 92 ++++++++++++++++++++++++
2 files changed, 112 insertions(+), 3 deletions(-)
create mode 100644 offload/test/offloading/target-no-loop.c
diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
index 3fccd3a291d37..973f091057250 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
@@ -701,6 +701,18 @@ static bool supportsSPMDExecutionMode(ASTContext &Ctx,
"Unknown programming model for OpenMP directive on NVPTX target.");
}
+/// Check whether a target kernel can be promoted to a "no-loop" SPMD kernel,
+/// mirroring Flang's MLIR promotion path.
+static bool canPromoteToNoLoop(const LangOptions &LangOpts,
+ const OMPExecutableDirective &D) {
+ OpenMPDirectiveKind DKind = D.getDirectiveKind();
+ 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>();
+}
+
void CGOpenMPRuntimeGPU::emitNonSPMDKernel(const OMPExecutableDirective &D,
StringRef ParentName,
llvm::Function *&OutlinedFn,
@@ -746,9 +758,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(CGM.getLangOpts(), 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;
diff --git a/offload/test/offloading/target-no-loop.c b/offload/test/offloading/target-no-loop.c
new file mode 100644
index 0000000000000..25e973cef68c8
--- /dev/null
+++ b/offload/test/offloading/target-no-loop.c
@@ -0,0 +1,92 @@
+// 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 fdd03ae41dd5e624d67540b243e800906a7b02d8 Mon Sep 17 00:00:00 2001
From: Nicole Aschenbrenner <nicole.aschenbrenner at amd.com>
Date: Wed, 29 Jul 2026 07:16:57 -0500
Subject: [PATCH 2/3] [OpenMP][IRBuilder] Fix offload entry creation on
per-function finalize
Add an additional guard to prevent offload entry creation on per
function finalize from Clang causing duplicate entries. Restricts entry
creation to module level finalize.
Prevents asserting on missing offload entries from nested
CodeGenFunction finalizing before module completion. Required for
no-loop kernel codegen to emit a canonical loop nest inside a target
region.
---
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index fc3812502fb2e..8b44b93147cb8 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -1037,7 +1037,7 @@ void OpenMPIRBuilder::finalize(Function *Fn) {
"OMPIRBuilder finalization \n";
};
- if (!OffloadInfoManager.empty())
+ if (!Fn && !OffloadInfoManager.empty())
createOffloadEntriesAndInfoMetadata(ErrorReportFn);
if (Config.EmitLLVMUsedMetaInfo.value_or(false)) {
>From bf644538fc16f1bf76dc8e467e7deda593733539 Mon Sep 17 00:00:00 2001
From: Nicole Aschenbrenner <nicole.aschenbrenner at amd.com>
Date: Wed, 29 Jul 2026 07:43:05 -0500
Subject: [PATCH 3/3] [OpenMP][Clang] Emit loop-free kernels for no-loop target
regions
Generate no-loop kernels without a loop over the iteration space,
matching Flang's no-loop implementation.
Requires -fopenmp-enable-irbuilder to emit through OpenMPIRBuilder,
allowing shared loop codegen with Flang.
---
clang/lib/CodeGen/CGOpenMPRuntime.h | 6 +++
clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp | 25 ++++++------
clang/lib/CodeGen/CGOpenMPRuntimeGPU.h | 4 ++
clang/lib/CodeGen/CGStmtOpenMP.cpp | 39 +++++++++++++++++--
clang/test/OpenMP/irbuilder_target_no_loop.c | 41 ++++++++++++++++++++
5 files changed, 99 insertions(+), 16 deletions(-)
create mode 100644 clang/test/OpenMP/irbuilder_target_no_loop.c
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.h b/clang/lib/CodeGen/CGOpenMPRuntime.h
index a81d3830a8035..dad39eab848f6 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 973f091057250..ac1bc34e56721 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
@@ -701,18 +701,6 @@ static bool supportsSPMDExecutionMode(ASTContext &Ctx,
"Unknown programming model for OpenMP directive on NVPTX target.");
}
-/// Check whether a target kernel can be promoted to a "no-loop" SPMD kernel,
-/// mirroring Flang's MLIR promotion path.
-static bool canPromoteToNoLoop(const LangOptions &LangOpts,
- const OMPExecutableDirective &D) {
- OpenMPDirectiveKind DKind = D.getDirectiveKind();
- 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>();
-}
-
void CGOpenMPRuntimeGPU::emitNonSPMDKernel(const OMPExecutableDirective &D,
StringRef ParentName,
llvm::Function *&OutlinedFn,
@@ -758,7 +746,7 @@ void CGOpenMPRuntimeGPU::emitKernelInit(const OMPExecutableDirective &D,
CodeGenFunction &CGF,
EntryFunctionState &EST, bool IsSPMD) {
llvm::OpenMPIRBuilder::TargetKernelDefaultAttrs Attrs;
- if (IsSPMD && canPromoteToNoLoop(CGM.getLangOpts(), D))
+ if (IsSPMD && canPromoteToNoLoop(D))
Attrs.ExecFlags =
llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_SPMD_NO_LOOP;
else
@@ -1157,6 +1145,17 @@ 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>();
+}
+
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 3a7ee5456a9d2..d1f682780b9bd 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 18524a7c32b7a..90b7fd68086c4 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 CGM.getLangOpts().OpenMPIRBuilder && 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.
@@ -3707,6 +3716,28 @@ emitInnerParallelForWhenCombined(CodeGenFunction &CGF,
HasCancel = D->hasCancel();
}
CodeGenFunction::OMPCancelStackRAII CancelRegion(CGF, EKind, HasCancel);
+
+ CodeGenModule &CGM = CGF.CGM;
+ if (canEmitGPUNoLoopKernel(CGM, S)) {
+ const Stmt *Inner = S.getRawStmt();
+ llvm::CanonicalLoopInfo *CLI =
+ CGF.EmitOMPCollapsedCanonicalLoopNest(Inner, 1);
+ 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));
+ return;
+ }
+
CGF.EmitOMPWorksharingLoop(S, S.getPrevEnsureUpperBound(),
emitDistributeParallelForInnerBounds,
emitDistributeParallelForDispatchBounds);
@@ -6309,9 +6340,11 @@ void CodeGenFunction::EmitOMPDistributeLoop(const OMPLoopDirective &S,
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()) ||
+ canEmitGPUNoLoopKernel(CGM, S)) {
JumpDest LoopExit =
getJumpDestInCurrentScope(createBasicBlock("omp.loop.exit"));
CodeGenLoop(*this, S, LoopExit);
diff --git a/clang/test/OpenMP/irbuilder_target_no_loop.c b/clang/test/OpenMP/irbuilder_target_no_loop.c
new file mode 100644
index 0000000000000..1e8e7af71cbef
--- /dev/null
+++ b/clang/test/OpenMP/irbuilder_target_no_loop.c
@@ -0,0 +1,41 @@
+// 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: -fopenmp-enable-irbuilder -emit-llvm %s -o - | FileCheck %s \
+// RUN: --check-prefixes=CHECK,IRB \
+// RUN: --implicit-check-not=__kmpc_distribute_static_init \
+// RUN: --implicit-check-not=__kmpc_for_static_init
+
+// 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=CHECK,NOIRB \
+// RUN: --implicit-check-not=__kmpc_distribute_for_static_loop
+
+// 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;
+}
+
+// The third field is OMP_TGT_EXEC_MODE_SPMD_NO_LOOP (1 << 2 | 1 << 1).
+// CHECK: @{{.*}}no_loop{{.*}}_kernel_environment = weak_odr protected addrspace(1) constant %struct.KernelEnvironmentTy { %struct.ConfigurationEnvironmentTy { i8 0, i8 1, i8 6,
+
+// The trailing i8 is the NoLoop flag, forwarded to the device runtime as
+// one_iteration_per_thread.
+// IRB: call void @__kmpc_distribute_for_static_loop_4u({{.*}}, i32 0, i32 0, i8 1)
+
+// NOIRB: call void @__kmpc_distribute_static_init_4(
+// NOIRB: call void @__kmpc_for_static_init_4(
More information about the cfe-commits
mailing list