[clang] [CIR][OpenMP] Add support for combined target parallel directives (PR #207019)
Jan Leyonberg via cfe-commits
cfe-commits at lists.llvm.org
Wed Sep 30 15:47:38 PDT 2026
https://github.com/jsjodin updated https://github.com/llvm/llvm-project/pull/207019
>From a8ec7387caf37a1ae16ad431643f9033ee07064a 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/9] Initial implementaion
---
clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp | 309 ++++++++++--------
.../test/CIR/CodeGenOpenMP/target-parallel.c | 55 ++++
2 files changed, 232 insertions(+), 132 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 eca46da4658b3..0201e53077fa1 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
@@ -16,10 +16,171 @@
#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;
+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;
+}
+
+/// 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').
+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(cgf, cgm, builder, begin, clauses);
+ ce.emitIf(clauseOps, llvm::omp::Directive::OMPD_parallel);
+ ce.emitNumThreads(clauseOps);
+ ce.emitProcBind(clauseOps);
+ ce.emitNYI</*supported=*/OMPIfClause, OMPNumThreadsClause, OMPProcBindClause>(
+ /*nyi=*/OpenMPNYIClauseList<
+ OMPAllocateClause, OMPCopyinClause, OMPDefaultClause,
+ OMPFirstprivateClause, OMPPrivateClause, OMPReductionClause,
+ OMPSharedClause>{},
+ llvm::omp::Directive::OMPD_parallel);
+
+ auto parallelOp = mlir::omp::ParallelOp::create(builder, begin, clauseOps);
+
+ mlir::Block &block = parallelOp.getRegion().emplaceBlock();
+ mlir::OpBuilder::InsertionGuard guard(builder);
+ builder.setInsertionPointToEnd(&block);
+
+ CIRGenFunction::LexicalScope ls{cgf, 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");
+
+ 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");
@@ -32,47 +193,17 @@ CIRGenFunction::emitOMPErrorDirective(const OMPErrorDirective &s) {
}
mlir::LogicalResult
CIRGenFunction::emitOMPParallelDirective(const OMPParallelDirective &s) {
- mlir::LogicalResult res = mlir::success();
mlir::Location begin = getLoc(s.getBeginLoc());
mlir::Location end = getLoc(s.getEndLoc());
- mlir::omp::ParallelOperands clauseOps;
- OpenMPClauseEmitter ce(*this, getCIRGenModule(), builder, begin, s.clauses());
- ce.emitIf(clauseOps, llvm::omp::Directive::OMPD_parallel);
- ce.emitNumThreads(clauseOps);
- ce.emitProcBind(clauseOps);
- ce.emitNYI</*supported=*/OMPIfClause, OMPNumThreadsClause, OMPProcBindClause>(
- /*nyi=*/OpenMPNYIClauseList<OMPAllocateClause, OMPCopyinClause,
- OMPDefaultClause, OMPFirstprivateClause,
- OMPPrivateClause, OMPReductionClause,
- OMPSharedClause>{},
- llvm::omp::Directive::OMPD_parallel);
-
- auto parallelOp = mlir::omp::ParallelOp::create(builder, begin, clauseOps);
-
- {
- mlir::Block &block = parallelOp.getRegion().emplaceBlock();
- mlir::OpBuilder::InsertionGuard guardCase(builder);
- builder.setInsertionPointToEnd(&block);
-
- LexicalScope ls{*this, begin, builder.getInsertionBlock()};
-
- if (s.hasCancel())
- getCIRGenModule().errorNYI(s.getBeginLoc(),
- "OpenMP Parallel with Cancel");
- if (s.getTaskReductionRefExpr())
- getCIRGenModule().errorNYI(s.getBeginLoc(),
- "OpenMP Parallel with Task Reduction");
+ 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
@@ -225,107 +356,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) {
@@ -363,9 +402,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 3feaa1618dd4a7f4ad01b75ff7a3e3a68b257c7c 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/9] 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 0201e53077fa1..5afce0216b971 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').
@@ -116,81 +91,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());
@@ -356,6 +256,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 854b94ca38947ed447090aaf537a4256903e1032 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/9] 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 5afce0216b971..fa1d2cf7bdab4 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,
@@ -290,10 +287,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 96ffd2282a575be644c7514b49c4e76918ae69da 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/9] Address review feedback
---
clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h | 15 +-
.../CIRGenOpenMPConstructDecomposition.h | 317 ++++++++++++++++++
clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp | 173 +++++++---
.../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 5 +-
.../CIR/CodeGenOpenMP/not-yet-implemented.c | 15 +
clang/test/CIR/CodeGenOpenMP/parallel.c | 7 +-
.../test/CIR/CodeGenOpenMP/target-parallel.c | 11 +-
7 files changed, 475 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 6a731a9128986..a5dcbcb918c73 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h
@@ -55,10 +55,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.
@@ -67,12 +68,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 =
@@ -81,11 +84,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 fa1d2cf7bdab4..92d82779cd1fe 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,42 +34,63 @@ 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.emitIf(clauseOps, llvm::omp::Directive::OMPD_parallel);
ce.emitNumThreads(clauseOps);
ce.emitProcBind(clauseOps);
- ce.emitNYI</*supported=*/OMPIfClause, OMPNumThreadsClause, OMPProcBindClause>(
+ return ce.emitNYI</*supported=*/OMPIfClause, OMPNumThreadsClause,
+ OMPProcBindClause>(
/*nyi=*/OpenMPNYIClauseList<
OMPAllocateClause, OMPCopyinClause, OMPDefaultClause,
OMPFirstprivateClause, 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);
@@ -78,10 +100,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);
@@ -93,14 +119,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
@@ -287,23 +321,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,
@@ -311,6 +337,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);
@@ -319,6 +354,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)
@@ -351,10 +388,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) {
@@ -395,12 +442,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 0d715e6023d04..9941bf994d9ac 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
@@ -5975,8 +5975,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 93a17494ad605..9a1068974db30 100644
--- a/clang/test/CIR/CodeGenOpenMP/not-yet-implemented.c
+++ b/clang/test/CIR/CodeGenOpenMP/not-yet-implemented.c
@@ -8,4 +8,19 @@ void do_things() {
// expected-error at +1{{ClangIR code gen Not Yet Implemented: OpenMP OMPSingleDirective}}
#pragma omp single
{}
+
+ int 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 49c5e10db7183..6a9a349c59489 100644
--- a/clang/test/CIR/CodeGenOpenMP/parallel.c
+++ b/clang/test/CIR/CodeGenOpenMP/parallel.c
@@ -1,4 +1,4 @@
-// RUN: not %clang_cc1 -fopenmp -triple x86_64-linux-gnu -emit-cir -fclangir %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -triple x86_64-linux-gnu -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);
>From 0aaa8804a43b71935251b982f9d1fbbafec5156e Mon Sep 17 00:00:00 2001
From: Jan Leyonberg <jan_sjodin at yahoo.com>
Date: Wed, 2 Sep 2026 13:29:41 -0400
Subject: [PATCH 5/9] Address second round of review feedback
---
.../CIRGenOpenMPConstructDecomposition.h | 21 ++-
clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp | 157 ++++++++++--------
.../CIR/CodeGenOpenMP/not-yet-implemented.c | 6 +
3 files changed, 114 insertions(+), 70 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenMPConstructDecomposition.h b/clang/lib/CIR/CodeGen/CIRGenOpenMPConstructDecomposition.h
index 3c54a6293e769..ec8c012502a41 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenMPConstructDecomposition.h
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenMPConstructDecomposition.h
@@ -286,18 +286,26 @@ struct LeafWithClauses {
llvm::SmallVector<llvm::omp::Clause> synthesized;
};
-inline llvm::SmallVector<LeafWithClauses>
-decompose(unsigned openmpVersion, const OMPExecutableDirective &s) {
+using ConstructQueue = llvm::SmallVector<LeafWithClauses>;
+
+/// Given a potentially compound directive with a list of clauses that apply to
+/// it, break it up into individual leaf constructs each with the subset of
+/// applicable clauses (plus implicit clauses, if any). From that create a work
+/// queue, ordered outermost to innermost, where each work item corresponds to
+/// the leaf construct with its clauses. Implicit clauses are synthesized by the
+/// decomposition and have no Clang AST node, so they are listed separately.
+inline ConstructQueue buildConstructQueue(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::omp::Version(openmpVersion), helper, s.getDirectiveKind(),
llvm::ArrayRef<Clause>(input));
- llvm::SmallVector<LeafWithClauses> result;
+ ConstructQueue result;
for (const tomp::DirectiveWithClauses<Clause> &dwc : decomp.output) {
LeafWithClauses leaf;
leaf.id = dwc.id;
@@ -312,6 +320,11 @@ decompose(unsigned openmpVersion, const OMPExecutableDirective &s) {
return result;
}
+inline bool isLastItemInQueue(ConstructQueue::const_iterator item,
+ const ConstructQueue &queue) {
+ return std::next(item) == queue.end();
+}
+
} // 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 92d82779cd1fe..fde91689e9c59 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
@@ -34,36 +34,21 @@ CIRGenFunction::emitOMPErrorDirective(const OMPErrorDirective &s) {
return mlir::failure();
}
-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 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);
+/// Report \p item's synthesized clauses as not-yet-implemented: having no AST
+/// node, they cannot go through the clause emitters.
+static mlir::LogicalResult
+checkSynthesizedClauses(CIRGenFunction &cgf, const OMPExecutableDirective &s,
+ omp::ConstructQueue::const_iterator item) {
+ mlir::LogicalResult res = mlir::success();
+ for (llvm::omp::Clause synth : item->synthesized) {
+ cgf.getCIRGenModule().errorNYI(s.getSourceRange(),
+ (llvm::Twine("OpenMP synthesized '") +
+ llvm::omp::getOpenMPClauseName(synth) +
+ "' clause from construct decomposition")
+ .str());
+ res = mlir::failure();
}
- return result;
-}
-
-/// 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);
-
- 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();
+ return res;
}
static mlir::LogicalResult
@@ -86,13 +71,17 @@ emitParallelClauses(CIRGenFunction &cgf, CIRGenModule &cgm,
template <typename DirectiveTy>
static mlir::LogicalResult
-emitParallelOp(CIRGenFunction &cgf, const DirectiveTy &s, mlir::Location begin,
+emitParallelOp(CIRGenFunction &cgf, const DirectiveTy &s,
+ const omp::ConstructQueue &queue,
+ omp::ConstructQueue::const_iterator item, 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);
+ if (!omp::isLastItemInQueue(item, queue))
+ parallelOp.setCombined(true);
mlir::Block &block = parallelOp.getRegion().emplaceBlock();
mlir::OpBuilder::InsertionGuard guard(builder);
@@ -119,15 +108,21 @@ CIRGenFunction::emitOMPParallelDirective(const OMPParallelDirective &s) {
mlir::Location begin = getLoc(s.getBeginLoc());
mlir::Location end = getLoc(s.getEndLoc());
- llvm::SmallVector<const OMPClause *> clauses =
- getLeafClauses(*this, s, llvm::omp::OMPD_parallel);
+ omp::ConstructQueue queue =
+ omp::buildConstructQueue(getContext().getLangOpts().OpenMP, s);
+ omp::ConstructQueue::const_iterator item = queue.begin();
+
+ if (mlir::failed(checkSynthesizedClauses(*this, s, item)))
+ return mlir::failure();
+
mlir::omp::ParallelOperands clauseOps;
if (mlir::failed(emitParallelClauses(*this, getCIRGenModule(), builder, begin,
- clauses, clauseOps)))
+ item->clauses, clauseOps)))
return mlir::failure();
return emitParallelOp(
- *this, s, begin, end, clauseOps, [&]() -> mlir::LogicalResult {
+ *this, s, queue, item, 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
@@ -288,26 +283,30 @@ CIRGenFunction::emitOMPAtomicDirective(const OMPAtomicDirective &s) {
}
/// Check for unsupported implicit captures in a target region.
-static void
+static mlir::LogicalResult
emitOMPTargetImplicitCaptures(CIRGenFunction &cgf,
const OMPExecutableDirective &s,
llvm::ArrayRef<const VarDecl *> mapSyms) {
const CapturedStmt *cs = s.getCapturedStmt(llvm::omp::OMPD_target);
+ mlir::LogicalResult res = mlir::success();
for (const auto &capture : cs->captures()) {
if (capture.capturesThis()) {
cgf.getCIRGenModule().errorNYI(s.getBeginLoc(),
"OpenMP target capture of 'this' pointer");
+ res = mlir::failure();
continue;
}
if (capture.capturesVariableByCopy()) {
cgf.getCIRGenModule().errorNYI(s.getBeginLoc(),
"OpenMP target capture by copy");
+ res = mlir::failure();
continue;
}
if (capture.capturesVariableArrayType()) {
cgf.getCIRGenModule().errorNYI(
s.getBeginLoc(),
"OpenMP target capture of variable-length array type");
+ res = mlir::failure();
continue;
}
if (capture.capturesVariable()) {
@@ -317,8 +316,10 @@ emitOMPTargetImplicitCaptures(CIRGenFunction &cgf,
cgf.getCIRGenModule().errorNYI(s.getBeginLoc(),
"OpenMP target implicit by-ref capture");
+ res = mlir::failure();
}
}
+ return res;
}
static mlir::LogicalResult
@@ -341,20 +342,23 @@ emitTargetClauses(CIRGenFunction &cgf, CIRGenModule &cgm,
template <typename DirectiveTy>
static mlir::LogicalResult
-emitTargetOp(CIRGenFunction &cgf, const DirectiveTy &s, mlir::Location begin,
+emitTargetOp(CIRGenFunction &cgf, const DirectiveTy &s,
+ const omp::ConstructQueue &queue,
+ omp::ConstructQueue::const_iterator item, 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);
+ if (mlir::failed(emitOMPTargetImplicitCaptures(cgf, s, mapSyms)))
+ return mlir::failure();
// 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);
- if (leafIsCombined(cgf, s, llvm::omp::OMPD_target))
+ if (!omp::isLastItemInQueue(item, queue))
targetOp.setCombined(true);
mlir::Block &block = targetOp.getRegion().emplaceBlock();
@@ -388,20 +392,26 @@ CIRGenFunction::emitOMPTargetDirective(const OMPTargetDirective &s) {
mlir::Location begin = getLoc(s.getBeginLoc());
mlir::Location end = getLoc(s.getEndLoc());
- llvm::SmallVector<const OMPClause *> clauses =
- getLeafClauses(*this, s, llvm::omp::OMPD_target);
+ omp::ConstructQueue queue =
+ omp::buildConstructQueue(getContext().getLangOpts().OpenMP, s);
+ omp::ConstructQueue::const_iterator item = queue.begin();
+
+ if (mlir::failed(checkSynthesizedClauses(*this, s, item)))
+ return mlir::failure();
+
mlir::omp::TargetExtOperands clauseOps;
llvm::SmallVector<const VarDecl *> mapSyms;
if (mlir::failed(emitTargetClauses(*this, getCIRGenModule(), builder, begin,
- clauses, clauseOps, mapSyms)))
+ item->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);
- });
+ return emitTargetOp(*this, s, queue, item, 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) {
@@ -442,31 +452,46 @@ mlir::LogicalResult CIRGenFunction::emitOMPTargetParallelDirective(
mlir::Location begin = getLoc(s.getBeginLoc());
mlir::Location end = getLoc(s.getEndLoc());
- // 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);
+ omp::ConstructQueue queue =
+ omp::buildConstructQueue(getContext().getLangOpts().OpenMP, s);
+ omp::ConstructQueue::const_iterator targetItem = queue.begin();
+ assert(targetItem->id == llvm::omp::OMPD_target &&
+ "expected 'target' to be the outermost leaf");
+
+ if (mlir::failed(checkSynthesizedClauses(*this, s, targetItem)))
+ return mlir::failure();
+
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)))
+ targetItem->clauses, targetOps, mapSyms)))
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);
- });
+ *this, s, queue, targetItem, begin, end, targetOps, mapSyms,
+ [&]() -> mlir::LogicalResult {
+ omp::ConstructQueue::const_iterator parallelItem =
+ std::next(targetItem);
+ assert(parallelItem != queue.end() &&
+ parallelItem->id == llvm::omp::OMPD_parallel &&
+ "expected a 'parallel' leaf nested in 'target'");
+
+ if (mlir::failed(checkSynthesizedClauses(*this, s, parallelItem)))
+ return mlir::failure();
+
+ mlir::omp::ParallelOperands parallelOps;
+ if (mlir::failed(emitParallelClauses(*this, getCIRGenModule(), builder,
+ begin, parallelItem->clauses,
+ parallelOps)))
+ return mlir::failure();
+
+ return emitParallelOp(*this, s, queue, parallelItem, begin, end,
+ parallelOps, [&]() -> mlir::LogicalResult {
+ const CapturedStmt *cs =
+ s.getCapturedStmt(llvm::omp::OMPD_parallel);
+ return emitStmt(cs->getCapturedStmt(),
+ /*useCurrentScope=*/true);
+ });
});
}
mlir::LogicalResult CIRGenFunction::emitOMPTargetParallelForDirective(
diff --git a/clang/test/CIR/CodeGenOpenMP/not-yet-implemented.c b/clang/test/CIR/CodeGenOpenMP/not-yet-implemented.c
index 9a1068974db30..c4a97e2b76d05 100644
--- a/clang/test/CIR/CodeGenOpenMP/not-yet-implemented.c
+++ b/clang/test/CIR/CodeGenOpenMP/not-yet-implemented.c
@@ -23,4 +23,10 @@ void do_things() {
// expected-error at +1{{ClangIR code gen Not Yet Implemented: OpenMP TARGET 'private' clause}}
#pragma omp target private(i)
{}
+
+ // Decomposition can also synthesize a clause that the user did not write and
+ // that has no Clang AST node to drive its emitter.
+ // expected-error at +1{{ClangIR code gen Not Yet Implemented: OpenMP synthesized 'firstprivate' clause from construct decomposition}}
+#pragma omp target
+ { i = 1; }
}
>From 542bbe13c9fafdea6a4bc0884ef114f5ebf39bbd Mon Sep 17 00:00:00 2001
From: Jan Leyonberg <jan_sjodin at yahoo.com>
Date: Mon, 28 Sep 2026 16:31:01 -0400
Subject: [PATCH 6/9] Fix if clause
---
.../CIRGenOpenMPConstructDecomposition.h | 34 +++++++++++++++++--
1 file changed, 31 insertions(+), 3 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenMPConstructDecomposition.h b/clang/lib/CIR/CodeGen/CIRGenOpenMPConstructDecomposition.h
index ec8c012502a41..063b4af1c462a 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenMPConstructDecomposition.h
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenMPConstructDecomposition.h
@@ -288,6 +288,27 @@ struct LeafWithClauses {
using ConstructQueue = llvm::SmallVector<LeafWithClauses>;
+/// The decomposition algorithm resolves a directive-name-modifier on an `if`
+/// clause (e.g. `if(parallel: cond)`) by synthesizing a fresh clause with the
+/// modifier stripped -- reusing the same condition expression -- and
+/// assigning it to the leaf named by the modifier. That synthesized clause
+/// has no Clang AST node of its own (\c original is null), but it shares its
+/// condition expression's identity with the clause it was derived from. Use
+/// that to recover the original \c OMPIfClause so a resolved `if` clause is
+/// routed to its leaf like any other clause, instead of being reported as an
+/// unimplemented synthesized clause.
+inline const clang::OMPClause *
+findOriginalIfClause(const OMPExecutableDirective &s, const Clause &c) {
+ const auto &ifClause = std::get<tomp::clause::IfT<TypeTy, IdTy, ExprTy>>(c.u);
+ const clang::Expr *condition = std::get<ExprTy>(ifClause.t);
+ for (const clang::OMPClause *oc : s.clauses()) {
+ const auto *ic = llvm::dyn_cast<clang::OMPIfClause>(oc);
+ if (ic && ic->getCondition() == condition)
+ return ic;
+ }
+ return nullptr;
+}
+
/// Given a potentially compound directive with a list of clauses that apply to
/// it, break it up into individual leaf constructs each with the subset of
/// applicable clauses (plus implicit clauses, if any). From that create a work
@@ -310,10 +331,17 @@ inline ConstructQueue buildConstructQueue(unsigned openmpVersion,
LeafWithClauses leaf;
leaf.id = dwc.id;
for (const Clause &c : dwc.clauses) {
- if (c.original)
+ if (c.original) {
leaf.clauses.push_back(c.original);
- else
- leaf.synthesized.push_back(c.id);
+ continue;
+ }
+ if (c.id == llvm::omp::OMPC_if) {
+ if (const clang::OMPClause *orig = findOriginalIfClause(s, c)) {
+ leaf.clauses.push_back(orig);
+ continue;
+ }
+ }
+ leaf.synthesized.push_back(c.id);
}
result.push_back(std::move(leaf));
}
>From 3d582c50173ebfc4b0db300a086c3a5bb85923a9 Mon Sep 17 00:00:00 2001
From: Jan Leyonberg <jan_sjodin at yahoo.com>
Date: Mon, 28 Sep 2026 16:51:37 -0400
Subject: [PATCH 7/9] Fix review comments
---
clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp | 16 ++++++++--------
1 file changed, 8 insertions(+), 8 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
index fde91689e9c59..e109d3e108404 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
@@ -89,14 +89,16 @@ emitParallelOp(CIRGenFunction &cgf, const DirectiveTy &s,
CIRGenFunction::LexicalScope ls{cgf, begin, builder.getInsertionBlock()};
+ // hasCancel() reports a nested cancel/cancellation point directive in the
+ // body, not a clause on `parallel`, so it can't be an emitParallelClauses()
+ // NYI check.
if (s.hasCancel()) {
cgm.errorNYI(s.getBeginLoc(), "OpenMP Parallel with Cancel");
return mlir::failure();
}
- if (s.getTaskReductionRefExpr()) {
- cgm.errorNYI(s.getBeginLoc(), "OpenMP Parallel with Task Reduction");
- return mlir::failure();
- }
+ // Only set for reduction(task: ...), already rejected as NYI above.
+ assert(!s.getTaskReductionRefExpr() &&
+ "reduction(task: ...) should already be rejected as NYI");
mlir::LogicalResult res = emitBody();
mlir::omp::TerminatorOp::create(builder, end);
@@ -123,10 +125,8 @@ CIRGenFunction::emitOMPParallelDirective(const OMPParallelDirective &s) {
return emitParallelOp(
*this, s, queue, item, 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.
+ // emitStmt() rejects CapturedStmt directly; the parent construct
+ // must unwrap it, so emit the inner statement instead.
const CapturedStmt *cs = s.getCapturedStmt(llvm::omp::OMPD_parallel);
return emitStmt(cs->getCapturedStmt(), /*useCurrentScope=*/true);
});
>From 3b6c739a49ce99048f75515d3a1a1d13d0c443d3 Mon Sep 17 00:00:00 2001
From: Jan Leyonberg <jan_sjodin at yahoo.com>
Date: Mon, 28 Sep 2026 16:56:19 -0400
Subject: [PATCH 8/9] fix formatting
---
clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp | 8 ++++----
1 file changed, 4 insertions(+), 4 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
index e109d3e108404..9d937f5ed67f9 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp
@@ -62,10 +62,10 @@ emitParallelClauses(CIRGenFunction &cgf, CIRGenModule &cgm,
ce.emitProcBind(clauseOps);
return ce.emitNYI</*supported=*/OMPIfClause, OMPNumThreadsClause,
OMPProcBindClause>(
- /*nyi=*/OpenMPNYIClauseList<
- OMPAllocateClause, OMPCopyinClause, OMPDefaultClause,
- OMPFirstprivateClause, OMPPrivateClause, OMPReductionClause,
- OMPSharedClause>{},
+ /*nyi=*/OpenMPNYIClauseList<OMPAllocateClause, OMPCopyinClause,
+ OMPDefaultClause, OMPFirstprivateClause,
+ OMPPrivateClause, OMPReductionClause,
+ OMPSharedClause>{},
llvm::omp::Directive::OMPD_parallel);
}
>From 83ca70aeec728b4c51bf2e9651f4d83b068e573a Mon Sep 17 00:00:00 2001
From: Jan Leyonberg <jan_sjodin at yahoo.com>
Date: Wed, 30 Sep 2026 18:45:12 -0400
Subject: [PATCH 9/9] Adapt to isClauseAllowedOnDirective decomposition helper
requirement
Upstream's ConstructDecompositionT now requires the Helper type to
implement isClauseAllowedOnDirective (llvm/llvm-project#227418),
replacing the free-function call it used to make directly. Add the
same delegating implementation flang uses.
Co-authored-by: Cursor <cursoragent at cursor.com>
---
clang/lib/CIR/CodeGen/CIRGenOpenMPConstructDecomposition.h | 7 +++++++
1 file changed, 7 insertions(+)
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenMPConstructDecomposition.h b/clang/lib/CIR/CodeGen/CIRGenOpenMPConstructDecomposition.h
index 063b4af1c462a..38c749e9eb253 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenMPConstructDecomposition.h
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenMPConstructDecomposition.h
@@ -278,6 +278,13 @@ struct DecompositionHelper {
}
/// CIR does not lower loop directives yet, so there is no iteration variable.
std::optional<Object> getLoopIterVar() const { return std::nullopt; }
+ /// Defer to the spec table; CIR does not override allowed clauses per
+ /// directive.
+ bool isClauseAllowedOnDirective(llvm::omp::Clause clauseId,
+ llvm::omp::Directive dirId,
+ llvm::omp::Version version) const {
+ return llvm::omp::isAllowedClauseForDirective(dirId, clauseId, version);
+ }
};
struct LeafWithClauses {
More information about the cfe-commits
mailing list