[clang] [CIR][OpenMP] Add support for combined target parallel directives (PR #207019)
Jan Leyonberg via cfe-commits
cfe-commits at lists.llvm.org
Sun Aug 30 17:15:51 PDT 2026
https://github.com/jsjodin updated https://github.com/llvm/llvm-project/pull/207019
>From 7d01b9ab51338895b0bed1d296d7c936d76debd8 Mon Sep 17 00:00:00 2001
From: Jan Leyonberg <jan_sjodin at yahoo.com>
Date: Wed, 1 Jul 2026 08:20:40 -0400
Subject: [PATCH 1/4] Initial implementaion
---
clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp | 303 ++++++++++--------
.../test/CIR/CodeGenOpenMP/target-parallel.c | 55 ++++
2 files changed, 229 insertions(+), 129 deletions(-)
create mode 100644 clang/test/CIR/CodeGenOpenMP/target-parallel.c
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
index a42735391629f..567141de05e27 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
@@ -16,28 +16,78 @@
#include "mlir/Dialect/OpenMP/OpenMPDialect.h"
#include "clang/AST/OpenMPClause.h"
#include "clang/AST/StmtOpenMP.h"
+#include "llvm/ADT/STLFunctionalExtras.h"
+#include "llvm/Frontend/OpenMP/OMP.h"
#include "llvm/Frontend/OpenMP/OMPConstants.h"
using namespace clang;
using namespace clang::CIRGen;
-mlir::LogicalResult
-CIRGenFunction::emitOMPScopeDirective(const OMPScopeDirective &s) {
- getCIRGenModule().errorNYI(s.getSourceRange(), "OpenMP OMPScopeDirective");
- return mlir::failure();
+namespace {
+
+/// Returns the subset of \p s's clauses that are allowed on the given leaf
+/// directive.
+static llvm::SmallVector<const OMPClause *>
+getLeafClauses(CIRGenFunction &cgf, const OMPExecutableDirective &s,
+ llvm::omp::Directive leaf) {
+ unsigned version = cgf.getContext().getLangOpts().OpenMP;
+ llvm::SmallVector<const OMPClause *> result;
+ for (const OMPClause *c : s.clauses())
+ if (llvm::omp::isAllowedClauseForDirective(leaf, c->getClauseKind(),
+ version))
+ result.push_back(c);
+ return result;
}
-mlir::LogicalResult
-CIRGenFunction::emitOMPErrorDirective(const OMPErrorDirective &s) {
- getCIRGenModule().errorNYI(s.getSourceRange(), "OpenMP OMPErrorDirective");
- return mlir::failure();
+
+/// Check for unsupported implicit captures in a target region.
+static void
+emitOMPTargetImplicitCaptures(CIRGenFunction &cgf,
+ const OMPExecutableDirective &s,
+ llvm::ArrayRef<const VarDecl *> mapSyms) {
+ const CapturedStmt *cs = s.getCapturedStmt(llvm::omp::OMPD_target);
+ for (const auto &capture : cs->captures()) {
+ if (capture.capturesThis()) {
+ cgf.getCIRGenModule().errorNYI(s.getBeginLoc(),
+ "OpenMP target capture of 'this' pointer");
+ continue;
+ }
+ if (capture.capturesVariableByCopy()) {
+ cgf.getCIRGenModule().errorNYI(s.getBeginLoc(),
+ "OpenMP target capture by copy");
+ continue;
+ }
+ if (capture.capturesVariableArrayType()) {
+ cgf.getCIRGenModule().errorNYI(
+ s.getBeginLoc(),
+ "OpenMP target capture of variable-length array type");
+ continue;
+ }
+ if (capture.capturesVariable()) {
+ const VarDecl *vd = capture.getCapturedVar();
+ if (llvm::is_contained(mapSyms, vd))
+ continue;
+
+ cgf.getCIRGenModule().errorNYI(s.getBeginLoc(),
+ "OpenMP target implicit by-ref capture");
+ }
+ }
}
-mlir::LogicalResult
-CIRGenFunction::emitOMPParallelDirective(const OMPParallelDirective &s) {
- mlir::LogicalResult res = mlir::success();
- mlir::Location begin = getLoc(s.getBeginLoc());
- mlir::Location end = getLoc(s.getEndLoc());
+
+/// Create an omp.parallel op for the parallel leaf of \p s and emit \p emitBody
+/// inside its region. Works for both the standalone 'parallel' directive and
+/// combined directives that contain a parallel leaf (e.g. 'target parallel').
+template <typename DirectiveTy>
+static mlir::LogicalResult
+emitParallelOp(CIRGenFunction &cgf, const DirectiveTy &s, mlir::Location begin,
+ mlir::Location end,
+ llvm::function_ref<mlir::LogicalResult()> emitBody) {
+ CIRGenBuilderTy &builder = cgf.getBuilder();
+ CIRGenModule &cgm = cgf.getCIRGenModule();
+
+ llvm::SmallVector<const OMPClause *> clauses =
+ getLeafClauses(cgf, s, llvm::omp::OMPD_parallel);
mlir::omp::ParallelOperands clauseOps;
- OpenMPClauseEmitter ce(*this, getCIRGenModule(), builder, begin, s.clauses());
+ OpenMPClauseEmitter ce(cgf, cgm, builder, begin, clauses);
ce.emitProcBind(clauseOps);
ce.emitNYI</*supported=*/OMPProcBindClause>(
/*nyi=*/OpenMPNYIClauseList<
@@ -48,29 +98,110 @@ CIRGenFunction::emitOMPParallelDirective(const OMPParallelDirective &s) {
auto parallelOp = mlir::omp::ParallelOp::create(builder, begin, clauseOps);
- {
- mlir::Block &block = parallelOp.getRegion().emplaceBlock();
- mlir::OpBuilder::InsertionGuard guardCase(builder);
- builder.setInsertionPointToEnd(&block);
+ mlir::Block &block = parallelOp.getRegion().emplaceBlock();
+ mlir::OpBuilder::InsertionGuard guard(builder);
+ builder.setInsertionPointToEnd(&block);
+
+ CIRGenFunction::LexicalScope ls{cgf, begin, builder.getInsertionBlock()};
- LexicalScope ls{*this, begin, builder.getInsertionBlock()};
+ if (s.hasCancel())
+ cgm.errorNYI(s.getBeginLoc(), "OpenMP Parallel with Cancel");
+ if (s.getTaskReductionRefExpr())
+ cgm.errorNYI(s.getBeginLoc(), "OpenMP Parallel with Task Reduction");
- if (s.hasCancel())
- getCIRGenModule().errorNYI(s.getBeginLoc(),
- "OpenMP Parallel with Cancel");
- if (s.getTaskReductionRefExpr())
- getCIRGenModule().errorNYI(s.getBeginLoc(),
- "OpenMP Parallel with Task Reduction");
+ mlir::LogicalResult res = emitBody();
+ mlir::omp::TerminatorOp::create(builder, end);
+ return res;
+}
+
+/// Create an omp.target op for the target leaf of \p s and emit \p emitBody
+/// inside its region, remapping mapped variables to the target op's block
+/// arguments. Works for both the standalone 'target' directive and combined
+/// directives that contain a target leaf (e.g. 'target parallel').
+template <typename DirectiveTy>
+static mlir::LogicalResult
+emitTargetOp(CIRGenFunction &cgf, const DirectiveTy &s, mlir::Location begin,
+ mlir::Location end,
+ llvm::function_ref<mlir::LogicalResult()> emitBody) {
+ CIRGenBuilderTy &builder = cgf.getBuilder();
+ CIRGenModule &cgm = cgf.getCIRGenModule();
+
+ llvm::SmallVector<const OMPClause *> clauses =
+ getLeafClauses(cgf, s, llvm::omp::OMPD_target);
+
+ mlir::omp::TargetExtOperands clauseOps;
+ llvm::SmallVector<const VarDecl *> mapSyms;
+
+ OpenMPClauseEmitter ce(cgf, cgm, builder, begin, clauses);
+ ce.emitMap(clauseOps, &mapSyms);
+ ce.emitNYI</*supported=*/OMPMapClause>(
+ /*nyi=*/OpenMPNYIClauseList<
+ OMPAllocateClause, OMPDefaultClause, OMPDefaultmapClause,
+ OMPDependClause, OMPDeviceClause, OMPFirstprivateClause,
+ OMPHasDeviceAddrClause, OMPIfClause, OMPInReductionClause,
+ OMPIsDevicePtrClause, OMPNowaitClause, OMPPrivateClause,
+ OMPThreadLimitClause, OMPUsesAllocatorsClause, OMPXBareClause>{},
+ llvm::omp::Directive::OMPD_target);
+
+ emitOMPTargetImplicitCaptures(cgf, s, mapSyms);
+
+ // Use generic for now.
+ clauseOps.kernelType = mlir::omp::TargetExecModeAttr::get(
+ &cgf.getMLIRContext(), mlir::omp::TargetExecMode::generic);
+
+ auto targetOp = mlir::omp::TargetOp::create(builder, begin, clauseOps);
+
+ mlir::Block &block = targetOp.getRegion().emplaceBlock();
+ for (mlir::Value mapVar : clauseOps.mapVars)
+ block.addArgument(mapVar.getType(), begin);
+
+ mlir::OpBuilder::InsertionGuard guard(builder);
+ builder.setInsertionPointToEnd(&block);
+
+ CIRGenFunction::LexicalScope ls{cgf, begin, builder.getInsertionBlock()};
+
+ llvm::SmallVector<std::pair<const VarDecl *, Address>> savedAddrs;
+ for (auto [idx, vd] : llvm::enumerate(mapSyms)) {
+ Address origAddr = cgf.getAddrOfLocalVar(vd);
+ savedAddrs.push_back({vd, origAddr});
+ mlir::Value blockArg = block.getArgument(idx);
+ cgf.replaceAddrOfLocalVar(vd, Address(blockArg, origAddr.getAlignment()));
+ }
+
+ mlir::LogicalResult res = emitBody();
+ mlir::omp::TerminatorOp::create(builder, end);
+
+ for (auto &[vd, addr] : savedAddrs)
+ cgf.replaceAddrOfLocalVar(vd, addr);
+
+ return res;
+}
+
+} // anonymous namespace
+
+mlir::LogicalResult
+CIRGenFunction::emitOMPScopeDirective(const OMPScopeDirective &s) {
+ getCIRGenModule().errorNYI(s.getSourceRange(), "OpenMP OMPScopeDirective");
+ return mlir::failure();
+}
+mlir::LogicalResult
+CIRGenFunction::emitOMPErrorDirective(const OMPErrorDirective &s) {
+ getCIRGenModule().errorNYI(s.getSourceRange(), "OpenMP OMPErrorDirective");
+ return mlir::failure();
+}
+mlir::LogicalResult
+CIRGenFunction::emitOMPParallelDirective(const OMPParallelDirective &s) {
+ mlir::Location begin = getLoc(s.getBeginLoc());
+ mlir::Location end = getLoc(s.getEndLoc());
+
+ return emitParallelOp(*this, s, begin, end, [&]() -> mlir::LogicalResult {
// Don't lower the captured statement directly since this will be
// special-cased depending on the kind of OpenMP directive that is the
// parent, also the non-OpenMP context captured statements lowering does
// not apply directly.
const CapturedStmt *cs = s.getCapturedStmt(llvm::omp::OMPD_parallel);
- const Stmt *bodyStmt = cs->getCapturedStmt();
- res = emitStmt(bodyStmt, /*useCurrentScope=*/true);
- mlir::omp::TerminatorOp::create(builder, end);
- }
- return res;
+ return emitStmt(cs->getCapturedStmt(), /*useCurrentScope=*/true);
+ });
}
mlir::LogicalResult
@@ -223,107 +354,15 @@ CIRGenFunction::emitOMPAtomicDirective(const OMPAtomicDirective &s) {
return mlir::failure();
}
-/// Check for unsupported implicit captures in a target region.
-static void
-emitOMPTargetImplicitCaptures(CIRGenFunction &cgf, const OMPTargetDirective &s,
- llvm::ArrayRef<const VarDecl *> mapSyms) {
- const CapturedStmt *cs = s.getCapturedStmt(llvm::omp::OMPD_target);
- for (const auto &capture : cs->captures()) {
- if (capture.capturesThis()) {
- cgf.getCIRGenModule().errorNYI(s.getBeginLoc(),
- "OpenMP target capture of 'this' pointer");
- continue;
- }
- if (capture.capturesVariableByCopy()) {
- cgf.getCIRGenModule().errorNYI(s.getBeginLoc(),
- "OpenMP target capture by copy");
- continue;
- }
- if (capture.capturesVariableArrayType()) {
- cgf.getCIRGenModule().errorNYI(
- s.getBeginLoc(),
- "OpenMP target capture of variable-length array type");
- continue;
- }
- if (capture.capturesVariable()) {
- const VarDecl *vd = capture.getCapturedVar();
- if (llvm::is_contained(mapSyms, vd))
- continue;
-
- cgf.getCIRGenModule().errorNYI(s.getBeginLoc(),
- "OpenMP target implicit by-ref capture");
- }
- }
-}
-
-/// Emit the body of an omp.target region, remapping mapped variables to the
-/// block arguments of the target op's region.
-static mlir::LogicalResult
-emitOMPTargetBody(CIRGenFunction &cgf, const OMPTargetDirective &s,
- mlir::omp::TargetOp targetOp,
- llvm::ArrayRef<mlir::Value> mapVars,
- llvm::ArrayRef<const VarDecl *> mappedVarDecls,
- mlir::Location begin, mlir::Location end) {
- mlir::Block &block = targetOp.getRegion().emplaceBlock();
-
- for (mlir::Value mapVar : mapVars)
- block.addArgument(mapVar.getType(), begin);
-
- mlir::OpBuilder::InsertionGuard guard(cgf.getBuilder());
- cgf.getBuilder().setInsertionPointToEnd(&block);
-
- CIRGenFunction::LexicalScope ls{cgf, begin,
- cgf.getBuilder().getInsertionBlock()};
-
- llvm::SmallVector<std::pair<const VarDecl *, Address>> savedAddrs;
- for (auto [idx, vd] : llvm::enumerate(mappedVarDecls)) {
- Address origAddr = cgf.getAddrOfLocalVar(vd);
- savedAddrs.push_back({vd, origAddr});
- mlir::Value blockArg = block.getArgument(idx);
- cgf.replaceAddrOfLocalVar(vd, Address(blockArg, origAddr.getAlignment()));
- }
-
- const CapturedStmt *cs = s.getCapturedStmt(llvm::omp::OMPD_target);
- mlir::LogicalResult res =
- cgf.emitStmt(cs->getCapturedStmt(), /*useCurrentScope=*/true);
-
- mlir::omp::TerminatorOp::create(cgf.getBuilder(), end);
-
- for (auto &[vd, addr] : savedAddrs)
- cgf.replaceAddrOfLocalVar(vd, addr);
-
- return res;
-}
-
mlir::LogicalResult
CIRGenFunction::emitOMPTargetDirective(const OMPTargetDirective &s) {
mlir::Location begin = getLoc(s.getBeginLoc());
mlir::Location end = getLoc(s.getEndLoc());
- mlir::omp::TargetExtOperands clauseOps;
- llvm::SmallVector<const VarDecl *> mapSyms;
-
- OpenMPClauseEmitter ce(*this, getCIRGenModule(), builder, begin, s.clauses());
- ce.emitMap(clauseOps, &mapSyms);
- ce.emitNYI</*supported=*/OMPMapClause>(
- /*nyi=*/OpenMPNYIClauseList<
- OMPAllocateClause, OMPDefaultClause, OMPDefaultmapClause,
- OMPDependClause, OMPDeviceClause, OMPFirstprivateClause,
- OMPHasDeviceAddrClause, OMPIfClause, OMPInReductionClause,
- OMPIsDevicePtrClause, OMPNowaitClause, OMPPrivateClause,
- OMPThreadLimitClause, OMPUsesAllocatorsClause, OMPXBareClause>{},
- llvm::omp::Directive::OMPD_target);
-
- emitOMPTargetImplicitCaptures(*this, s, mapSyms);
-
- // Use generic for now.
- clauseOps.kernelType = mlir::omp::TargetExecModeAttr::get(
- &getMLIRContext(), mlir::omp::TargetExecMode::generic);
-
- auto targetOp = mlir::omp::TargetOp::create(builder, begin, clauseOps);
-
- return emitOMPTargetBody(*this, s, targetOp, clauseOps.mapVars, mapSyms,
- begin, end);
+ return emitTargetOp(*this, s, begin, end, [&]() -> mlir::LogicalResult {
+ const CapturedStmt *cs = s.getCapturedStmt(llvm::omp::OMPD_target);
+ return emitStmt(cs->getCapturedStmt(), /*useCurrentScope=*/true);
+ });
}
mlir::LogicalResult
CIRGenFunction::emitOMPTeamsDirective(const OMPTeamsDirective &s) {
@@ -361,9 +400,15 @@ mlir::LogicalResult CIRGenFunction::emitOMPTargetExitDataDirective(
}
mlir::LogicalResult CIRGenFunction::emitOMPTargetParallelDirective(
const OMPTargetParallelDirective &s) {
- getCIRGenModule().errorNYI(s.getSourceRange(),
- "OpenMP OMPTargetParallelDirective");
- return mlir::failure();
+ mlir::Location begin = getLoc(s.getBeginLoc());
+ mlir::Location end = getLoc(s.getEndLoc());
+
+ return emitTargetOp(*this, s, begin, end, [&]() -> mlir::LogicalResult {
+ return emitParallelOp(*this, s, begin, end, [&]() -> mlir::LogicalResult {
+ const CapturedStmt *cs = s.getCapturedStmt(llvm::omp::OMPD_parallel);
+ return emitStmt(cs->getCapturedStmt(), /*useCurrentScope=*/true);
+ });
+ });
}
mlir::LogicalResult CIRGenFunction::emitOMPTargetParallelForDirective(
const OMPTargetParallelForDirective &s) {
diff --git a/clang/test/CIR/CodeGenOpenMP/target-parallel.c b/clang/test/CIR/CodeGenOpenMP/target-parallel.c
new file mode 100644
index 0000000000000..eb9f1ea343eff
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenMP/target-parallel.c
@@ -0,0 +1,55 @@
+// Host compilation (x86 host, AMDGPU offload target).
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -emit-cir -fclangir %s -o - \
+// RUN: | FileCheck %s --check-prefix=CIR-HOST
+
+// Device compilation (AMDGPU): allocas live in the private address space.
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fopenmp -fopenmp-is-target-device \
+// RUN: -emit-cir -fclangir %s -o - \
+// RUN: | FileCheck %s --check-prefix=CIR-DEVICE
+
+void use(int);
+
+// The combined 'target parallel' directive lowers to an omp.parallel nested
+// inside an omp.target, identical to the equivalent nesting of the separate
+// 'target' and 'parallel' directives.
+void target_parallel(int x) {
+ // CIR-HOST: cir.func{{.*}}@target_parallel
+ // CIR-HOST: %[[MAP:.*]] = omp.map.info {{.*}} map_clauses(tofrom) {{.*}} {name = "x"}
+ // CIR-HOST: omp.target kernel_type(generic) map_entries(%[[MAP]] -> %[[ARG:.*]] : !cir.ptr<!s32i>) {
+ // CIR-HOST: omp.parallel {
+ // CIR-HOST: %[[LOAD:.*]] = cir.load align(4) %[[ARG]]
+ // CIR-HOST: cir.call @use(%[[LOAD]])
+ // CIR-HOST: omp.terminator
+ // CIR-HOST: }
+ // CIR-HOST: omp.terminator
+ // CIR-HOST: }
+
+ // CIR-DEVICE: cir.func{{.*}}@target_parallel
+ // CIR-DEVICE: omp.target kernel_type(generic) {{.*}} {
+ // CIR-DEVICE: omp.parallel {
+ // CIR-DEVICE: cir.call @use
+ // CIR-DEVICE: omp.terminator
+ // CIR-DEVICE: }
+ // CIR-DEVICE: omp.terminator
+ // CIR-DEVICE: }
+#pragma omp target parallel map(tofrom : x)
+ {
+ use(x);
+ }
+}
+
+// 'target parallel' routes the proc_bind clause to the parallel leaf and the
+// map clause to the target leaf.
+void target_parallel_proc_bind(int x) {
+ // CIR-HOST: cir.func{{.*}}@target_parallel_proc_bind
+ // CIR-HOST: omp.target kernel_type(generic) map_entries({{.*}}) {
+ // CIR-HOST: omp.parallel proc_bind(spread) {
+ // CIR-HOST: omp.terminator
+ // CIR-HOST: }
+ // CIR-HOST: omp.terminator
+ // CIR-HOST: }
+#pragma omp target parallel proc_bind(spread) map(tofrom : x)
+ {
+ use(x);
+ }
+}
>From e961d5583cc003566c99b664302e77749718c03e Mon Sep 17 00:00:00 2001
From: Jan Leyonberg <jan_sjodin at yahoo.com>
Date: Wed, 1 Jul 2026 12:19:22 -0400
Subject: [PATCH 2/4] Move code around to reduce the diff.x
---
clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp | 217 ++++++++++-----------
1 file changed, 107 insertions(+), 110 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
index 567141de05e27..fd324335ec234 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
@@ -22,7 +22,16 @@
using namespace clang;
using namespace clang::CIRGen;
-namespace {
+mlir::LogicalResult
+CIRGenFunction::emitOMPScopeDirective(const OMPScopeDirective &s) {
+ getCIRGenModule().errorNYI(s.getSourceRange(), "OpenMP OMPScopeDirective");
+ return mlir::failure();
+}
+mlir::LogicalResult
+CIRGenFunction::emitOMPErrorDirective(const OMPErrorDirective &s) {
+ getCIRGenModule().errorNYI(s.getSourceRange(), "OpenMP OMPErrorDirective");
+ return mlir::failure();
+}
/// Returns the subset of \p s's clauses that are allowed on the given leaf
/// directive.
@@ -38,40 +47,6 @@ getLeafClauses(CIRGenFunction &cgf, const OMPExecutableDirective &s,
return result;
}
-/// Check for unsupported implicit captures in a target region.
-static void
-emitOMPTargetImplicitCaptures(CIRGenFunction &cgf,
- const OMPExecutableDirective &s,
- llvm::ArrayRef<const VarDecl *> mapSyms) {
- const CapturedStmt *cs = s.getCapturedStmt(llvm::omp::OMPD_target);
- for (const auto &capture : cs->captures()) {
- if (capture.capturesThis()) {
- cgf.getCIRGenModule().errorNYI(s.getBeginLoc(),
- "OpenMP target capture of 'this' pointer");
- continue;
- }
- if (capture.capturesVariableByCopy()) {
- cgf.getCIRGenModule().errorNYI(s.getBeginLoc(),
- "OpenMP target capture by copy");
- continue;
- }
- if (capture.capturesVariableArrayType()) {
- cgf.getCIRGenModule().errorNYI(
- s.getBeginLoc(),
- "OpenMP target capture of variable-length array type");
- continue;
- }
- if (capture.capturesVariable()) {
- const VarDecl *vd = capture.getCapturedVar();
- if (llvm::is_contained(mapSyms, vd))
- continue;
-
- cgf.getCIRGenModule().errorNYI(s.getBeginLoc(),
- "OpenMP target implicit by-ref capture");
- }
- }
-}
-
/// Create an omp.parallel op for the parallel leaf of \p s and emit \p emitBody
/// inside its region. Works for both the standalone 'parallel' directive and
/// combined directives that contain a parallel leaf (e.g. 'target parallel').
@@ -114,81 +89,6 @@ emitParallelOp(CIRGenFunction &cgf, const DirectiveTy &s, mlir::Location begin,
return res;
}
-/// Create an omp.target op for the target leaf of \p s and emit \p emitBody
-/// inside its region, remapping mapped variables to the target op's block
-/// arguments. Works for both the standalone 'target' directive and combined
-/// directives that contain a target leaf (e.g. 'target parallel').
-template <typename DirectiveTy>
-static mlir::LogicalResult
-emitTargetOp(CIRGenFunction &cgf, const DirectiveTy &s, mlir::Location begin,
- mlir::Location end,
- llvm::function_ref<mlir::LogicalResult()> emitBody) {
- CIRGenBuilderTy &builder = cgf.getBuilder();
- CIRGenModule &cgm = cgf.getCIRGenModule();
-
- llvm::SmallVector<const OMPClause *> clauses =
- getLeafClauses(cgf, s, llvm::omp::OMPD_target);
-
- mlir::omp::TargetExtOperands clauseOps;
- llvm::SmallVector<const VarDecl *> mapSyms;
-
- OpenMPClauseEmitter ce(cgf, cgm, builder, begin, clauses);
- ce.emitMap(clauseOps, &mapSyms);
- ce.emitNYI</*supported=*/OMPMapClause>(
- /*nyi=*/OpenMPNYIClauseList<
- OMPAllocateClause, OMPDefaultClause, OMPDefaultmapClause,
- OMPDependClause, OMPDeviceClause, OMPFirstprivateClause,
- OMPHasDeviceAddrClause, OMPIfClause, OMPInReductionClause,
- OMPIsDevicePtrClause, OMPNowaitClause, OMPPrivateClause,
- OMPThreadLimitClause, OMPUsesAllocatorsClause, OMPXBareClause>{},
- llvm::omp::Directive::OMPD_target);
-
- emitOMPTargetImplicitCaptures(cgf, s, mapSyms);
-
- // Use generic for now.
- clauseOps.kernelType = mlir::omp::TargetExecModeAttr::get(
- &cgf.getMLIRContext(), mlir::omp::TargetExecMode::generic);
-
- auto targetOp = mlir::omp::TargetOp::create(builder, begin, clauseOps);
-
- mlir::Block &block = targetOp.getRegion().emplaceBlock();
- for (mlir::Value mapVar : clauseOps.mapVars)
- block.addArgument(mapVar.getType(), begin);
-
- mlir::OpBuilder::InsertionGuard guard(builder);
- builder.setInsertionPointToEnd(&block);
-
- CIRGenFunction::LexicalScope ls{cgf, begin, builder.getInsertionBlock()};
-
- llvm::SmallVector<std::pair<const VarDecl *, Address>> savedAddrs;
- for (auto [idx, vd] : llvm::enumerate(mapSyms)) {
- Address origAddr = cgf.getAddrOfLocalVar(vd);
- savedAddrs.push_back({vd, origAddr});
- mlir::Value blockArg = block.getArgument(idx);
- cgf.replaceAddrOfLocalVar(vd, Address(blockArg, origAddr.getAlignment()));
- }
-
- mlir::LogicalResult res = emitBody();
- mlir::omp::TerminatorOp::create(builder, end);
-
- for (auto &[vd, addr] : savedAddrs)
- cgf.replaceAddrOfLocalVar(vd, addr);
-
- return res;
-}
-
-} // anonymous namespace
-
-mlir::LogicalResult
-CIRGenFunction::emitOMPScopeDirective(const OMPScopeDirective &s) {
- getCIRGenModule().errorNYI(s.getSourceRange(), "OpenMP OMPScopeDirective");
- return mlir::failure();
-}
-mlir::LogicalResult
-CIRGenFunction::emitOMPErrorDirective(const OMPErrorDirective &s) {
- getCIRGenModule().errorNYI(s.getSourceRange(), "OpenMP OMPErrorDirective");
- return mlir::failure();
-}
mlir::LogicalResult
CIRGenFunction::emitOMPParallelDirective(const OMPParallelDirective &s) {
mlir::Location begin = getLoc(s.getBeginLoc());
@@ -354,6 +254,103 @@ CIRGenFunction::emitOMPAtomicDirective(const OMPAtomicDirective &s) {
return mlir::failure();
}
+/// Check for unsupported implicit captures in a target region.
+static void
+emitOMPTargetImplicitCaptures(CIRGenFunction &cgf,
+ const OMPExecutableDirective &s,
+ llvm::ArrayRef<const VarDecl *> mapSyms) {
+ const CapturedStmt *cs = s.getCapturedStmt(llvm::omp::OMPD_target);
+ for (const auto &capture : cs->captures()) {
+ if (capture.capturesThis()) {
+ cgf.getCIRGenModule().errorNYI(s.getBeginLoc(),
+ "OpenMP target capture of 'this' pointer");
+ continue;
+ }
+ if (capture.capturesVariableByCopy()) {
+ cgf.getCIRGenModule().errorNYI(s.getBeginLoc(),
+ "OpenMP target capture by copy");
+ continue;
+ }
+ if (capture.capturesVariableArrayType()) {
+ cgf.getCIRGenModule().errorNYI(
+ s.getBeginLoc(),
+ "OpenMP target capture of variable-length array type");
+ continue;
+ }
+ if (capture.capturesVariable()) {
+ const VarDecl *vd = capture.getCapturedVar();
+ if (llvm::is_contained(mapSyms, vd))
+ continue;
+
+ cgf.getCIRGenModule().errorNYI(s.getBeginLoc(),
+ "OpenMP target implicit by-ref capture");
+ }
+ }
+}
+
+/// Create an omp.target op for the target leaf of \p s and emit \p emitBody
+/// inside its region, remapping mapped variables to the target op's block
+/// arguments. Works for both the standalone 'target' directive and combined
+/// directives that contain a target leaf (e.g. 'target parallel').
+template <typename DirectiveTy>
+static mlir::LogicalResult
+emitTargetOp(CIRGenFunction &cgf, const DirectiveTy &s, mlir::Location begin,
+ mlir::Location end,
+ llvm::function_ref<mlir::LogicalResult()> emitBody) {
+ CIRGenBuilderTy &builder = cgf.getBuilder();
+ CIRGenModule &cgm = cgf.getCIRGenModule();
+
+ llvm::SmallVector<const OMPClause *> clauses =
+ getLeafClauses(cgf, s, llvm::omp::OMPD_target);
+
+ mlir::omp::TargetExtOperands clauseOps;
+ llvm::SmallVector<const VarDecl *> mapSyms;
+
+ OpenMPClauseEmitter ce(cgf, cgm, builder, begin, clauses);
+ ce.emitMap(clauseOps, &mapSyms);
+ ce.emitNYI</*supported=*/OMPMapClause>(
+ /*nyi=*/OpenMPNYIClauseList<
+ OMPAllocateClause, OMPDefaultClause, OMPDefaultmapClause,
+ OMPDependClause, OMPDeviceClause, OMPFirstprivateClause,
+ OMPHasDeviceAddrClause, OMPIfClause, OMPInReductionClause,
+ OMPIsDevicePtrClause, OMPNowaitClause, OMPPrivateClause,
+ OMPThreadLimitClause, OMPUsesAllocatorsClause, OMPXBareClause>{},
+ llvm::omp::Directive::OMPD_target);
+
+ emitOMPTargetImplicitCaptures(cgf, s, mapSyms);
+
+ // Use generic for now.
+ clauseOps.kernelType = mlir::omp::TargetExecModeAttr::get(
+ &cgf.getMLIRContext(), mlir::omp::TargetExecMode::generic);
+
+ auto targetOp = mlir::omp::TargetOp::create(builder, begin, clauseOps);
+
+ mlir::Block &block = targetOp.getRegion().emplaceBlock();
+ for (mlir::Value mapVar : clauseOps.mapVars)
+ block.addArgument(mapVar.getType(), begin);
+
+ mlir::OpBuilder::InsertionGuard guard(builder);
+ builder.setInsertionPointToEnd(&block);
+
+ CIRGenFunction::LexicalScope ls{cgf, begin, builder.getInsertionBlock()};
+
+ llvm::SmallVector<std::pair<const VarDecl *, Address>> savedAddrs;
+ for (auto [idx, vd] : llvm::enumerate(mapSyms)) {
+ Address origAddr = cgf.getAddrOfLocalVar(vd);
+ savedAddrs.push_back({vd, origAddr});
+ mlir::Value blockArg = block.getArgument(idx);
+ cgf.replaceAddrOfLocalVar(vd, Address(blockArg, origAddr.getAlignment()));
+ }
+
+ mlir::LogicalResult res = emitBody();
+ mlir::omp::TerminatorOp::create(builder, end);
+
+ for (auto &[vd, addr] : savedAddrs)
+ cgf.replaceAddrOfLocalVar(vd, addr);
+
+ return res;
+}
+
mlir::LogicalResult
CIRGenFunction::emitOMPTargetDirective(const OMPTargetDirective &s) {
mlir::Location begin = getLoc(s.getBeginLoc());
>From fd303d70981ab002e155581dde8065800926da5f Mon Sep 17 00:00:00 2001
From: Jan Leyonberg <jan_sjodin at yahoo.com>
Date: Wed, 1 Jul 2026 12:22:55 -0400
Subject: [PATCH 3/4] Remove usless comments
---
clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp | 7 -------
1 file changed, 7 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
index fd324335ec234..76e5be3614a1d 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
@@ -47,9 +47,6 @@ getLeafClauses(CIRGenFunction &cgf, const OMPExecutableDirective &s,
return result;
}
-/// Create an omp.parallel op for the parallel leaf of \p s and emit \p emitBody
-/// inside its region. Works for both the standalone 'parallel' directive and
-/// combined directives that contain a parallel leaf (e.g. 'target parallel').
template <typename DirectiveTy>
static mlir::LogicalResult
emitParallelOp(CIRGenFunction &cgf, const DirectiveTy &s, mlir::Location begin,
@@ -288,10 +285,6 @@ emitOMPTargetImplicitCaptures(CIRGenFunction &cgf,
}
}
-/// Create an omp.target op for the target leaf of \p s and emit \p emitBody
-/// inside its region, remapping mapped variables to the target op's block
-/// arguments. Works for both the standalone 'target' directive and combined
-/// directives that contain a target leaf (e.g. 'target parallel').
template <typename DirectiveTy>
static mlir::LogicalResult
emitTargetOp(CIRGenFunction &cgf, const DirectiveTy &s, mlir::Location begin,
>From c7966ffc2d7d84f8c44b739aded649f4f8ecf9b4 Mon Sep 17 00:00:00 2001
From: Jan Leyonberg <jan_sjodin at yahoo.com>
Date: Fri, 28 Aug 2026 12:53:49 -0400
Subject: [PATCH 4/4] Address review feedback
---
clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h | 15 +-
.../CIRGenOpenMPConstructDecomposition.h | 317 ++++++++++++++++++
clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp | 172 +++++++---
.../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 5 +-
.../CIR/CodeGenOpenMP/not-yet-implemented.c | 14 +
clang/test/CIR/CodeGenOpenMP/parallel.c | 7 +-
.../test/CIR/CodeGenOpenMP/target-parallel.c | 11 +-
7 files changed, 473 insertions(+), 68 deletions(-)
create mode 100644 clang/lib/CIR/CodeGen/CIRGenOpenMPConstructDecomposition.h
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h b/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h
index 54c7366b1d769..2c1141e5fb541 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h
@@ -50,10 +50,11 @@ class OpenMPClauseEmitter {
/// Verify the clauses of a directive to make sure all legal cases are either
/// implemented or give a NYI error. The \p SupportedClauses and \p
/// NYIClauses type lists must be disjoint and cover all clauses eligible for
- /// the directive being processed.
+ /// the directive being processed. Returns failure if any not-yet-implemented
+ /// clause was present (and an error was emitted), success otherwise.
template <typename... SupportedClauses, typename... NYIClauses>
- void emitNYI(OpenMPNYIClauseList<NYIClauses...> nyi,
- llvm::omp::Directive directive) const;
+ mlir::LogicalResult emitNYI(OpenMPNYIClauseList<NYIClauses...> nyi,
+ llvm::omp::Directive directive) const;
private:
/// True if T is the same type as any of Ts.
@@ -62,12 +63,14 @@ class OpenMPClauseEmitter {
};
template <typename... SupportedClauses, typename... NYIClauses>
-void OpenMPClauseEmitter::emitNYI(OpenMPNYIClauseList<NYIClauses...>,
- llvm::omp::Directive directive) const {
+mlir::LogicalResult
+OpenMPClauseEmitter::emitNYI(OpenMPNYIClauseList<NYIClauses...>,
+ llvm::omp::Directive directive) const {
static_assert(
(!isAnyOf<NYIClauses, SupportedClauses...> && ...),
"the supported and not-yet-implemented clause lists must be disjoint");
+ mlir::LogicalResult result = mlir::success();
for (const OMPClause *c : clauses) {
if (isa<NYIClauses...>(c)) {
std::string msg =
@@ -76,11 +79,13 @@ void OpenMPClauseEmitter::emitNYI(OpenMPNYIClauseList<NYIClauses...>,
llvm::omp::getOpenMPClauseName(c->getClauseKind()) + "' clause")
.str();
cgm.errorNYI(c->getBeginLoc(), msg);
+ result = mlir::failure();
} else if (!isa<SupportedClauses...>(c)) {
// Unknown/illegal clause encountered.
llvm_unreachable("unexpected OpenMP clause");
}
}
+ return result;
}
} // namespace clang::CIRGen
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenMPConstructDecomposition.h b/clang/lib/CIR/CodeGen/CIRGenOpenMPConstructDecomposition.h
new file mode 100644
index 0000000000000..3c54a6293e769
--- /dev/null
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenMPConstructDecomposition.h
@@ -0,0 +1,317 @@
+//===--- CIRGenOpenMPConstructDecomposition.h -----------------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_CLANG_LIB_CIR_CODEGEN_CIRGENOPENMPCONSTRUCTDECOMPOSITION_H
+#define LLVM_CLANG_LIB_CIR_CODEGEN_CIRGENOPENMPCONSTRUCTDECOMPOSITION_H
+
+#include "clang/AST/Decl.h"
+#include "clang/AST/Expr.h"
+#include "clang/AST/OpenMPClause.h"
+#include "clang/AST/StmtOpenMP.h"
+#include "clang/Basic/OpenMPKinds.h"
+#include "llvm/ADT/ArrayRef.h"
+#include "llvm/ADT/SmallVector.h"
+#include "llvm/Frontend/OpenMP/ClauseT.h"
+#include "llvm/Frontend/OpenMP/ConstructDecompositionT.h"
+#include "llvm/Frontend/OpenMP/OMP.h"
+#include "llvm/Support/Casting.h"
+
+#include <cassert>
+#include <optional>
+#include <utility>
+
+namespace clang::CIRGen::omp {
+
+// tomp type parameters for Clang: objects are identified by their canonical
+// declaration, expressions and types are the corresponding AST nodes.
+using TypeTy = const clang::Type *;
+using IdTy = const clang::ValueDecl *;
+using ExprTy = const clang::Expr *;
+
+} // namespace clang::CIRGen::omp
+
+// The decomposition operates on tomp::ObjectT<IdTy, ExprTy>; provide the
+// specialization for our identity/expression types.
+namespace tomp::type {
+template <>
+struct ObjectT<clang::CIRGen::omp::IdTy, clang::CIRGen::omp::ExprTy> {
+ using IdTy = clang::CIRGen::omp::IdTy;
+ using ExprTy = clang::CIRGen::omp::ExprTy;
+
+ IdTy id() const { return identity; }
+ const std::optional<ExprTy> &ref() const { return designator; }
+
+ bool operator<(const ObjectT &other) const {
+ return identity < other.identity;
+ }
+
+ IdTy identity = nullptr;
+ std::optional<ExprTy> designator;
+};
+} // namespace tomp::type
+
+namespace clang::CIRGen::omp {
+
+using Object = tomp::ObjectT<IdTy, ExprTy>;
+using ObjectList = tomp::ObjectListT<IdTy, ExprTy>;
+using ClauseBase = tomp::ClauseT<TypeTy, IdTy, ExprTy>;
+
+/// A tomp clause that remembers the Clang AST clause it came from, so the
+/// existing emitters can emit it after the decomposition assigns it to a leaf.
+/// Synthesized clauses have no AST counterpart and leave `original` null.
+struct Clause : public ClauseBase {
+ Clause() = default;
+ Clause(ClauseBase &&base) : ClauseBase(std::move(base)) {}
+ const clang::OMPClause *original = nullptr;
+};
+
+inline const clang::ValueDecl *getBaseValueDecl(const clang::Expr *e) {
+ e = e->IgnoreParenImpCasts();
+ for (;;) {
+ if (const auto *ase = llvm::dyn_cast<clang::ArraySubscriptExpr>(e)) {
+ e = ase->getBase()->IgnoreParenImpCasts();
+ continue;
+ }
+ if (const auto *ase = llvm::dyn_cast<clang::ArraySectionExpr>(e)) {
+ e = ase->getBase()->IgnoreParenImpCasts();
+ continue;
+ }
+ break;
+ }
+ if (const auto *dre = llvm::dyn_cast<clang::DeclRefExpr>(e))
+ return llvm::cast<clang::ValueDecl>(dre->getDecl()->getCanonicalDecl());
+ if (const auto *me = llvm::dyn_cast<clang::MemberExpr>(e))
+ return llvm::cast<clang::ValueDecl>(
+ me->getMemberDecl()->getCanonicalDecl());
+ return nullptr;
+}
+
+/// Build a tomp Object whose identity is the base variable's canonical decl.
+inline Object makeObject(const clang::Expr *e) {
+ return Object{getBaseValueDecl(e), e};
+}
+
+/// Build the tomp object list from a Clang var-list clause.
+template <typename ClangClause>
+inline ObjectList makeObjects(const ClangClause &c) {
+ ObjectList list;
+ for (const clang::Expr *e : c.getVarRefs())
+ list.push_back(makeObject(e));
+ return list;
+}
+
+/// Wrap a tomp clause payload into a Clause, remembering its AST origin.
+template <typename Specific>
+inline Clause makeClause(llvm::omp::Clause id, Specific &&specific,
+ const clang::OMPClause &original) {
+ Clause c{ClauseBase{id, std::forward<Specific>(specific)}};
+ c.original = &original;
+ return c;
+}
+
+/// Clause kinds that need a dedicated conversion: they either have a specific
+/// applyClause() overload (so the payload type selects it) or their contents
+/// feed the algorithm. Guards the generic fallback in makeGeneric.
+inline bool needsSpecificHandling(llvm::omp::Clause kind) {
+ switch (kind) {
+ case llvm::omp::OMPC_allocate:
+ case llvm::omp::OMPC_collapse:
+ case llvm::omp::OMPC_default:
+ case llvm::omp::OMPC_dyn_groupprivate:
+ case llvm::omp::OMPC_firstprivate:
+ case llvm::omp::OMPC_if:
+ case llvm::omp::OMPC_lastprivate:
+ case llvm::omp::OMPC_linear:
+ case llvm::omp::OMPC_map:
+ case llvm::omp::OMPC_nowait:
+ case llvm::omp::OMPC_ompx_attribute:
+ case llvm::omp::OMPC_ompx_bare:
+ case llvm::omp::OMPC_order:
+ case llvm::omp::OMPC_private:
+ case llvm::omp::OMPC_reduction:
+ case llvm::omp::OMPC_shared:
+ case llvm::omp::OMPC_thread_limit:
+ return true;
+ default:
+ return false;
+ }
+}
+
+/// Clause kinds CIR is able to emit today.
+inline bool isEmittableClause(llvm::omp::Clause kind) {
+ switch (kind) {
+ case llvm::omp::OMPC_map:
+ case llvm::omp::OMPC_proc_bind:
+ return true;
+ default:
+ return false;
+ }
+}
+
+/// Represent a clause by kind only, using an inert empty payload that routes
+/// through the algorithm's generic applyClause() path (which reads just the
+/// clause id). Valid for any clause with no specific applyClause() overload.
+inline Clause makeGeneric(llvm::omp::Clause id, const clang::OMPClause &orig) {
+ assert((!isEmittableClause(id) || !needsSpecificHandling(id)) &&
+ "CIR-emittable clause needs specific decomposition handling");
+ return makeClause(id, tomp::clause::ThreadsT<TypeTy, IdTy, ExprTy>{}, orig);
+}
+
+/// Convert a single Clang clause to its tomp representation. Every kind is
+/// handled; contents are populated only where the algorithm reads them.
+inline Clause convertClause(const clang::OMPClause &c) {
+ namespace tc = tomp::clause;
+ const llvm::omp::Clause kind = c.getClauseKind();
+ switch (kind) {
+ // Clauses whose contents the algorithm inspects.
+ case llvm::omp::OMPC_map: {
+ tc::MapT<TypeTy, IdTy, ExprTy> m{
+ {/*MapType=*/std::nullopt, /*MapTypeModifiers=*/std::nullopt,
+ /*AttachModifier=*/std::nullopt, /*RefModifier=*/std::nullopt,
+ /*Mappers=*/std::nullopt, /*Iterator=*/std::nullopt,
+ /*LocatorList=*/makeObjects(llvm::cast<clang::OMPMapClause>(c))}};
+ return makeClause(kind, std::move(m), c);
+ }
+ case llvm::omp::OMPC_firstprivate:
+ return makeClause(
+ kind,
+ tc::FirstprivateT<TypeTy, IdTy, ExprTy>{
+ /*List=*/makeObjects(llvm::cast<clang::OMPFirstprivateClause>(c))},
+ c);
+ case llvm::omp::OMPC_private:
+ return makeClause(kind,
+ tc::PrivateT<TypeTy, IdTy, ExprTy>{/*List=*/makeObjects(
+ llvm::cast<clang::OMPPrivateClause>(c))},
+ c);
+ case llvm::omp::OMPC_shared:
+ return makeClause(kind,
+ tc::SharedT<TypeTy, IdTy, ExprTy>{/*List=*/makeObjects(
+ llvm::cast<clang::OMPSharedClause>(c))},
+ c);
+ case llvm::omp::OMPC_lastprivate:
+ return makeClause(
+ kind,
+ tc::LastprivateT<TypeTy, IdTy, ExprTy>{
+ {/*LastprivateModifier=*/std::nullopt,
+ /*List=*/makeObjects(llvm::cast<clang::OMPLastprivateClause>(c))}},
+ c);
+ case llvm::omp::OMPC_linear:
+ return makeClause(
+ kind,
+ tc::LinearT<TypeTy, IdTy, ExprTy>{
+ {/*StepComplexModifier=*/std::nullopt,
+ /*LinearModifier=*/std::nullopt,
+ /*List=*/makeObjects(llvm::cast<clang::OMPLinearClause>(c))}},
+ c);
+ case llvm::omp::OMPC_reduction:
+ return makeClause(
+ kind,
+ tc::ReductionT<TypeTy, IdTy, ExprTy>{
+ {/*ReductionModifier=*/std::nullopt, /*ReductionIdentifiers=*/{},
+ /*List=*/makeObjects(llvm::cast<clang::OMPReductionClause>(c))}},
+ c);
+ case llvm::omp::OMPC_if: {
+ const auto &ic = llvm::cast<clang::OMPIfClause>(c);
+ std::optional<llvm::omp::Directive> mod;
+ if (ic.getNameModifier() != llvm::omp::OMPD_unknown)
+ mod = ic.getNameModifier();
+ return makeClause(
+ kind,
+ tc::IfT<TypeTy, IdTy, ExprTy>{{/*DirectiveNameModifier=*/mod,
+ /*IfExpression=*/ic.getCondition()}},
+ c);
+ }
+ // Clauses with a specific applyClause() overload but no contents the
+ // algorithm reads: carry the correct payload type so dispatch selects it.
+ case llvm::omp::OMPC_allocate:
+ return makeClause(kind,
+ tc::AllocateT<TypeTy, IdTy, ExprTy>{
+ {std::nullopt, std::nullopt, /*List=*/{}}},
+ c);
+ case llvm::omp::OMPC_collapse:
+ return makeClause(kind, tc::CollapseT<TypeTy, IdTy, ExprTy>{/*N=*/nullptr},
+ c);
+ case llvm::omp::OMPC_default:
+ return makeClause(
+ kind,
+ tc::DefaultT<TypeTy, IdTy, ExprTy>{
+ tc::DefaultT<TypeTy, IdTy, ExprTy>::DataSharingAttribute::Shared},
+ c);
+ case llvm::omp::OMPC_dyn_groupprivate:
+ return makeClause(kind,
+ tc::DynGroupprivateT<TypeTy, IdTy, ExprTy>{
+ {std::nullopt, std::nullopt, /*Size=*/nullptr}},
+ c);
+ case llvm::omp::OMPC_nowait:
+ return makeClause(kind, tc::NowaitT<TypeTy, IdTy, ExprTy>{}, c);
+ case llvm::omp::OMPC_ompx_attribute:
+ return makeClause(kind, tc::OmpxAttributeT<TypeTy, IdTy, ExprTy>{}, c);
+ case llvm::omp::OMPC_ompx_bare:
+ return makeClause(kind, tc::OmpxBareT<TypeTy, IdTy, ExprTy>{}, c);
+ case llvm::omp::OMPC_order:
+ return makeClause(
+ kind,
+ tc::OrderT<TypeTy, IdTy, ExprTy>{
+ {std::nullopt,
+ tc::OrderT<TypeTy, IdTy, ExprTy>::Ordering::Concurrent}},
+ c);
+ case llvm::omp::OMPC_thread_limit:
+ return makeClause(kind, tc::ThreadLimitT<TypeTy, IdTy, ExprTy>{/*List=*/{}},
+ c);
+ // Everything else routes by kind alone.
+ default:
+ return makeGeneric(kind, c);
+ }
+}
+
+/// Helper required by ConstructDecompositionT.
+struct DecompositionHelper {
+ /// Our object identities are already normalized to the base variable's decl,
+ /// so an object is its own base.
+ std::optional<Object> getBaseObject(const Object &object) const {
+ return object;
+ }
+ /// CIR does not lower loop directives yet, so there is no iteration variable.
+ std::optional<Object> getLoopIterVar() const { return std::nullopt; }
+};
+
+struct LeafWithClauses {
+ llvm::omp::Directive id = llvm::omp::Directive::OMPD_unknown;
+ llvm::SmallVector<const clang::OMPClause *> clauses;
+ llvm::SmallVector<llvm::omp::Clause> synthesized;
+};
+
+inline llvm::SmallVector<LeafWithClauses>
+decompose(unsigned openmpVersion, const OMPExecutableDirective &s) {
+ llvm::SmallVector<Clause> input;
+ for (const OMPClause *c : s.clauses())
+ input.push_back(convertClause(*c));
+
+ DecompositionHelper helper;
+ tomp::ConstructDecompositionT<Clause, DecompositionHelper> decomp(
+ openmpVersion, helper, s.getDirectiveKind(),
+ llvm::ArrayRef<Clause>(input));
+
+ llvm::SmallVector<LeafWithClauses> result;
+ for (const tomp::DirectiveWithClauses<Clause> &dwc : decomp.output) {
+ LeafWithClauses leaf;
+ leaf.id = dwc.id;
+ for (const Clause &c : dwc.clauses) {
+ if (c.original)
+ leaf.clauses.push_back(c.original);
+ else
+ leaf.synthesized.push_back(c.id);
+ }
+ result.push_back(std::move(leaf));
+ }
+ return result;
+}
+
+} // namespace clang::CIRGen::omp
+
+#endif // LLVM_CLANG_LIB_CIR_CODEGEN_CIRGENOPENMPCONSTRUCTDECOMPOSITION_H
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
index 76e5be3614a1d..69b77a8f34a9b 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
@@ -13,6 +13,7 @@
#include "CIRGenBuilder.h"
#include "CIRGenFunction.h"
#include "CIRGenOpenMPClause.h"
+#include "CIRGenOpenMPConstructDecomposition.h"
#include "mlir/Dialect/OpenMP/OpenMPDialect.h"
#include "clang/AST/OpenMPClause.h"
#include "clang/AST/StmtOpenMP.h"
@@ -33,40 +34,60 @@ CIRGenFunction::emitOMPErrorDirective(const OMPErrorDirective &s) {
return mlir::failure();
}
-/// Returns the subset of \p s's clauses that are allowed on the given leaf
-/// directive.
static llvm::SmallVector<const OMPClause *>
getLeafClauses(CIRGenFunction &cgf, const OMPExecutableDirective &s,
llvm::omp::Directive leaf) {
unsigned version = cgf.getContext().getLangOpts().OpenMP;
+ llvm::SmallVector<omp::LeafWithClauses> leaves = omp::decompose(version, s);
+
llvm::SmallVector<const OMPClause *> result;
- for (const OMPClause *c : s.clauses())
- if (llvm::omp::isAllowedClauseForDirective(leaf, c->getClauseKind(),
- version))
- result.push_back(c);
+ for (const omp::LeafWithClauses &l : leaves) {
+ if (l.id != leaf)
+ continue;
+ for (llvm::omp::Clause synth : l.synthesized)
+ cgf.getCIRGenModule().errorNYI(s.getSourceRange(),
+ (llvm::Twine("OpenMP synthesized '") +
+ llvm::omp::getOpenMPClauseName(synth) +
+ "' clause from construct decomposition")
+ .str());
+ llvm::append_range(result, l.clauses);
+ }
return result;
}
-template <typename DirectiveTy>
-static mlir::LogicalResult
-emitParallelOp(CIRGenFunction &cgf, const DirectiveTy &s, mlir::Location begin,
- mlir::Location end,
- llvm::function_ref<mlir::LogicalResult()> emitBody) {
- CIRGenBuilderTy &builder = cgf.getBuilder();
- CIRGenModule &cgm = cgf.getCIRGenModule();
+/// True when the op emitted for \p leaf must carry the omp.combined marker.
+static bool leafIsCombined(CIRGenFunction &cgf, const OMPExecutableDirective &s,
+ llvm::omp::Directive leaf) {
+ unsigned version = cgf.getContext().getLangOpts().OpenMP;
+ llvm::SmallVector<omp::LeafWithClauses> leaves = omp::decompose(version, s);
- llvm::SmallVector<const OMPClause *> clauses =
- getLeafClauses(cgf, s, llvm::omp::OMPD_parallel);
+ const auto *it = llvm::find_if(
+ leaves, [leaf](const omp::LeafWithClauses &l) { return l.id == leaf; });
+ return it != leaves.end() && std::next(it) != leaves.end();
+}
- mlir::omp::ParallelOperands clauseOps;
- OpenMPClauseEmitter ce(cgf, cgm, builder, begin, clauses);
+static mlir::LogicalResult
+emitParallelClauses(CIRGenFunction &cgf, CIRGenModule &cgm,
+ CIRGenBuilderTy &builder, mlir::Location loc,
+ llvm::ArrayRef<const OMPClause *> clauses,
+ mlir::omp::ParallelOperands &clauseOps) {
+ OpenMPClauseEmitter ce(cgf, cgm, builder, loc, clauses);
ce.emitProcBind(clauseOps);
- ce.emitNYI</*supported=*/OMPProcBindClause>(
+ return ce.emitNYI</*supported=*/OMPProcBindClause>(
/*nyi=*/OpenMPNYIClauseList<
OMPAllocateClause, OMPCopyinClause, OMPDefaultClause,
OMPFirstprivateClause, OMPIfClause, OMPNumThreadsClause,
OMPPrivateClause, OMPReductionClause, OMPSharedClause>{},
llvm::omp::Directive::OMPD_parallel);
+}
+
+template <typename DirectiveTy>
+static mlir::LogicalResult
+emitParallelOp(CIRGenFunction &cgf, const DirectiveTy &s, mlir::Location begin,
+ mlir::Location end, const mlir::omp::ParallelOperands &clauseOps,
+ llvm::function_ref<mlir::LogicalResult()> emitBody) {
+ CIRGenBuilderTy &builder = cgf.getBuilder();
+ CIRGenModule &cgm = cgf.getCIRGenModule();
auto parallelOp = mlir::omp::ParallelOp::create(builder, begin, clauseOps);
@@ -76,10 +97,14 @@ emitParallelOp(CIRGenFunction &cgf, const DirectiveTy &s, mlir::Location begin,
CIRGenFunction::LexicalScope ls{cgf, begin, builder.getInsertionBlock()};
- if (s.hasCancel())
+ if (s.hasCancel()) {
cgm.errorNYI(s.getBeginLoc(), "OpenMP Parallel with Cancel");
- if (s.getTaskReductionRefExpr())
+ return mlir::failure();
+ }
+ if (s.getTaskReductionRefExpr()) {
cgm.errorNYI(s.getBeginLoc(), "OpenMP Parallel with Task Reduction");
+ return mlir::failure();
+ }
mlir::LogicalResult res = emitBody();
mlir::omp::TerminatorOp::create(builder, end);
@@ -91,14 +116,22 @@ CIRGenFunction::emitOMPParallelDirective(const OMPParallelDirective &s) {
mlir::Location begin = getLoc(s.getBeginLoc());
mlir::Location end = getLoc(s.getEndLoc());
- return emitParallelOp(*this, s, begin, end, [&]() -> mlir::LogicalResult {
- // Don't lower the captured statement directly since this will be
- // special-cased depending on the kind of OpenMP directive that is the
- // parent, also the non-OpenMP context captured statements lowering does
- // not apply directly.
- const CapturedStmt *cs = s.getCapturedStmt(llvm::omp::OMPD_parallel);
- return emitStmt(cs->getCapturedStmt(), /*useCurrentScope=*/true);
- });
+ llvm::SmallVector<const OMPClause *> clauses =
+ getLeafClauses(*this, s, llvm::omp::OMPD_parallel);
+ mlir::omp::ParallelOperands clauseOps;
+ if (mlir::failed(emitParallelClauses(*this, getCIRGenModule(), builder, begin,
+ clauses, clauseOps)))
+ return mlir::failure();
+
+ return emitParallelOp(
+ *this, s, begin, end, clauseOps, [&]() -> mlir::LogicalResult {
+ // Don't lower the captured statement directly since this will be
+ // special-cased depending on the kind of OpenMP directive that is the
+ // parent, also the non-OpenMP context captured statements lowering does
+ // not apply directly.
+ const CapturedStmt *cs = s.getCapturedStmt(llvm::omp::OMPD_parallel);
+ return emitStmt(cs->getCapturedStmt(), /*useCurrentScope=*/true);
+ });
}
mlir::LogicalResult
@@ -285,23 +318,15 @@ emitOMPTargetImplicitCaptures(CIRGenFunction &cgf,
}
}
-template <typename DirectiveTy>
static mlir::LogicalResult
-emitTargetOp(CIRGenFunction &cgf, const DirectiveTy &s, mlir::Location begin,
- mlir::Location end,
- llvm::function_ref<mlir::LogicalResult()> emitBody) {
- CIRGenBuilderTy &builder = cgf.getBuilder();
- CIRGenModule &cgm = cgf.getCIRGenModule();
-
- llvm::SmallVector<const OMPClause *> clauses =
- getLeafClauses(cgf, s, llvm::omp::OMPD_target);
-
- mlir::omp::TargetExtOperands clauseOps;
- llvm::SmallVector<const VarDecl *> mapSyms;
-
- OpenMPClauseEmitter ce(cgf, cgm, builder, begin, clauses);
+emitTargetClauses(CIRGenFunction &cgf, CIRGenModule &cgm,
+ CIRGenBuilderTy &builder, mlir::Location loc,
+ llvm::ArrayRef<const OMPClause *> clauses,
+ mlir::omp::TargetExtOperands &clauseOps,
+ llvm::SmallVectorImpl<const VarDecl *> &mapSyms) {
+ OpenMPClauseEmitter ce(cgf, cgm, builder, loc, clauses);
ce.emitMap(clauseOps, &mapSyms);
- ce.emitNYI</*supported=*/OMPMapClause>(
+ return ce.emitNYI</*supported=*/OMPMapClause>(
/*nyi=*/OpenMPNYIClauseList<
OMPAllocateClause, OMPDefaultClause, OMPDefaultmapClause,
OMPDependClause, OMPDeviceClause, OMPFirstprivateClause,
@@ -309,6 +334,15 @@ emitTargetOp(CIRGenFunction &cgf, const DirectiveTy &s, mlir::Location begin,
OMPIsDevicePtrClause, OMPNowaitClause, OMPPrivateClause,
OMPThreadLimitClause, OMPUsesAllocatorsClause, OMPXBareClause>{},
llvm::omp::Directive::OMPD_target);
+}
+
+template <typename DirectiveTy>
+static mlir::LogicalResult
+emitTargetOp(CIRGenFunction &cgf, const DirectiveTy &s, mlir::Location begin,
+ mlir::Location end, mlir::omp::TargetExtOperands &clauseOps,
+ llvm::ArrayRef<const VarDecl *> mapSyms,
+ llvm::function_ref<mlir::LogicalResult()> emitBody) {
+ CIRGenBuilderTy &builder = cgf.getBuilder();
emitOMPTargetImplicitCaptures(cgf, s, mapSyms);
@@ -317,6 +351,8 @@ emitTargetOp(CIRGenFunction &cgf, const DirectiveTy &s, mlir::Location begin,
&cgf.getMLIRContext(), mlir::omp::TargetExecMode::generic);
auto targetOp = mlir::omp::TargetOp::create(builder, begin, clauseOps);
+ if (leafIsCombined(cgf, s, llvm::omp::OMPD_target))
+ targetOp.setCombined(true);
mlir::Block &block = targetOp.getRegion().emplaceBlock();
for (mlir::Value mapVar : clauseOps.mapVars)
@@ -349,10 +385,20 @@ CIRGenFunction::emitOMPTargetDirective(const OMPTargetDirective &s) {
mlir::Location begin = getLoc(s.getBeginLoc());
mlir::Location end = getLoc(s.getEndLoc());
- return emitTargetOp(*this, s, begin, end, [&]() -> mlir::LogicalResult {
- const CapturedStmt *cs = s.getCapturedStmt(llvm::omp::OMPD_target);
- return emitStmt(cs->getCapturedStmt(), /*useCurrentScope=*/true);
- });
+ llvm::SmallVector<const OMPClause *> clauses =
+ getLeafClauses(*this, s, llvm::omp::OMPD_target);
+ mlir::omp::TargetExtOperands clauseOps;
+ llvm::SmallVector<const VarDecl *> mapSyms;
+ if (mlir::failed(emitTargetClauses(*this, getCIRGenModule(), builder, begin,
+ clauses, clauseOps, mapSyms)))
+ return mlir::failure();
+
+ return emitTargetOp(
+ *this, s, begin, end, clauseOps, mapSyms, [&]() -> mlir::LogicalResult {
+ const CapturedStmt *cs = s.getCapturedStmt(llvm::omp::OMPD_target);
+ return emitStmt(cs->getCapturedStmt(),
+ /*useCurrentScope=*/true);
+ });
}
mlir::LogicalResult
CIRGenFunction::emitOMPTeamsDirective(const OMPTeamsDirective &s) {
@@ -393,12 +439,32 @@ mlir::LogicalResult CIRGenFunction::emitOMPTargetParallelDirective(
mlir::Location begin = getLoc(s.getBeginLoc());
mlir::Location end = getLoc(s.getEndLoc());
- return emitTargetOp(*this, s, begin, end, [&]() -> mlir::LogicalResult {
- return emitParallelOp(*this, s, begin, end, [&]() -> mlir::LogicalResult {
- const CapturedStmt *cs = s.getCapturedStmt(llvm::omp::OMPD_parallel);
- return emitStmt(cs->getCapturedStmt(), /*useCurrentScope=*/true);
- });
- });
+ // Split the clauses per leaf construct and evaluate them into their operand
+ // structures before creating the nested target/parallel ops.
+ llvm::SmallVector<const OMPClause *> targetClauses =
+ getLeafClauses(*this, s, llvm::omp::OMPD_target);
+ mlir::omp::TargetExtOperands targetOps;
+ llvm::SmallVector<const VarDecl *> mapSyms;
+ if (mlir::failed(emitTargetClauses(*this, getCIRGenModule(), builder, begin,
+ targetClauses, targetOps, mapSyms)))
+ return mlir::failure();
+
+ llvm::SmallVector<const OMPClause *> parallelClauses =
+ getLeafClauses(*this, s, llvm::omp::OMPD_parallel);
+ mlir::omp::ParallelOperands parallelOps;
+ if (mlir::failed(emitParallelClauses(*this, getCIRGenModule(), builder, begin,
+ parallelClauses, parallelOps)))
+ return mlir::failure();
+
+ return emitTargetOp(
+ *this, s, begin, end, targetOps, mapSyms, [&]() -> mlir::LogicalResult {
+ return emitParallelOp(
+ *this, s, begin, end, parallelOps, [&]() -> mlir::LogicalResult {
+ const CapturedStmt *cs =
+ s.getCapturedStmt(llvm::omp::OMPD_parallel);
+ return emitStmt(cs->getCapturedStmt(), /*useCurrentScope=*/true);
+ });
+ });
}
mlir::LogicalResult CIRGenFunction::emitOMPTargetParallelForDirective(
const OMPTargetParallelForDirective &s) {
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
index 9ee597b81df24..f53ff9a72a25a 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
@@ -5865,8 +5865,11 @@ void populateCIRToLLVMPasses(mlir::OpPassManager &pm, bool enableOpenMP) {
if (enableOpenMP)
pm.addPass(mlir::omp::createMarkDeclareTargetPass());
pm.addPass(createConvertCIRToLLVMPass());
- if (enableOpenMP)
+ if (enableOpenMP) {
pm.addPass(mlir::omp::createHostOpFilteringPass());
+ pm.nest<mlir::LLVM::LLVMFuncOp>().addPass(
+ mlir::omp::createStackToSharedPass());
+ }
}
std::unique_ptr<llvm::Module>
diff --git a/clang/test/CIR/CodeGenOpenMP/not-yet-implemented.c b/clang/test/CIR/CodeGenOpenMP/not-yet-implemented.c
index 29e280bf262d0..e3587ae0a71d2 100644
--- a/clang/test/CIR/CodeGenOpenMP/not-yet-implemented.c
+++ b/clang/test/CIR/CodeGenOpenMP/not-yet-implemented.c
@@ -13,4 +13,18 @@ void do_things() {
// expected-error at +1{{ClangIR code gen Not Yet Implemented: OpenMP PARALLEL 'if' clause}}
#pragma omp parallel if(i)
{}
+
+ // A leaf that reports a not-yet-implemented clause emits no op at all, rather
+ // than one that silently ignores the clause.
+ int a, b;
+ // expected-error at +2{{ClangIR code gen Not Yet Implemented: OpenMP PARALLEL 'shared' clause}}
+ // expected-error at +1{{ClangIR code gen Not Yet Implemented: OpenMP PARALLEL 'firstprivate' clause}}
+#pragma omp parallel shared(a) firstprivate(b)
+ {}
+
+ // A clause routed through construct decomposition but not yet emittable must
+ // still be diagnosed by the leaf emitter's NYI handling.
+ // expected-error at +1{{ClangIR code gen Not Yet Implemented: OpenMP TARGET 'private' clause}}
+#pragma omp target private(i)
+ {}
}
diff --git a/clang/test/CIR/CodeGenOpenMP/parallel.c b/clang/test/CIR/CodeGenOpenMP/parallel.c
index 36a48b38ee789..36e69557f93b4 100644
--- a/clang/test/CIR/CodeGenOpenMP/parallel.c
+++ b/clang/test/CIR/CodeGenOpenMP/parallel.c
@@ -1,4 +1,4 @@
-// RUN: not %clang_cc1 -fopenmp -emit-cir -fclangir %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -emit-cir -fclangir %s -o - | FileCheck %s
void before(int);
void during(int);
@@ -36,10 +36,7 @@ void parallel_with_operations() {
int a, b;
// CHECK-NEXT: cir.alloca "a"
// CHECK-NEXT: cir.alloca "b"
- // TODO(OMP): At the moment this results in 3 NYI diagnostics, 1 each for the
- // clauses + 1 for the CapturedStmt. When those are implemented, the check
- // lines will need updating.
-#pragma omp parallel shared(a) firstprivate(b)
+#pragma omp parallel
{
a = a + 1;
b = b + 1;
diff --git a/clang/test/CIR/CodeGenOpenMP/target-parallel.c b/clang/test/CIR/CodeGenOpenMP/target-parallel.c
index eb9f1ea343eff..a1832e7ffb3c2 100644
--- a/clang/test/CIR/CodeGenOpenMP/target-parallel.c
+++ b/clang/test/CIR/CodeGenOpenMP/target-parallel.c
@@ -1,3 +1,5 @@
+// REQUIRES: amdgpu-registered-target
+
// Host compilation (x86 host, AMDGPU offload target).
// RUN: %clang_cc1 -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -emit-cir -fclangir %s -o - \
// RUN: | FileCheck %s --check-prefix=CIR-HOST
@@ -14,7 +16,7 @@ void use(int);
// 'target' and 'parallel' directives.
void target_parallel(int x) {
// CIR-HOST: cir.func{{.*}}@target_parallel
- // CIR-HOST: %[[MAP:.*]] = omp.map.info {{.*}} map_clauses(tofrom) {{.*}} {name = "x"}
+ // CIR-HOST: %[[MAP:.*]] = omp.map.info {{.*}} map_clauses(tofrom) {{.*}} name("x")
// CIR-HOST: omp.target kernel_type(generic) map_entries(%[[MAP]] -> %[[ARG:.*]] : !cir.ptr<!s32i>) {
// CIR-HOST: omp.parallel {
// CIR-HOST: %[[LOAD:.*]] = cir.load align(4) %[[ARG]]
@@ -22,7 +24,8 @@ void target_parallel(int x) {
// CIR-HOST: omp.terminator
// CIR-HOST: }
// CIR-HOST: omp.terminator
- // CIR-HOST: }
+ // The target is the non-innermost leaf of the combined construct.
+ // CIR-HOST: } {omp.combined}
// CIR-DEVICE: cir.func{{.*}}@target_parallel
// CIR-DEVICE: omp.target kernel_type(generic) {{.*}} {
@@ -31,7 +34,7 @@ void target_parallel(int x) {
// CIR-DEVICE: omp.terminator
// CIR-DEVICE: }
// CIR-DEVICE: omp.terminator
- // CIR-DEVICE: }
+ // CIR-DEVICE: } {omp.combined}
#pragma omp target parallel map(tofrom : x)
{
use(x);
@@ -47,7 +50,7 @@ void target_parallel_proc_bind(int x) {
// CIR-HOST: omp.terminator
// CIR-HOST: }
// CIR-HOST: omp.terminator
- // CIR-HOST: }
+ // CIR-HOST: } {omp.combined}
#pragma omp target parallel proc_bind(spread) map(tofrom : x)
{
use(x);
More information about the cfe-commits
mailing list