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 <[email protected]> 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 <[email protected]> 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 <[email protected]> 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 <[email protected]> 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@+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@+2{{ClangIR code gen Not Yet Implemented: OpenMP PARALLEL 'shared' clause}} + // expected-error@+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@+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); _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
