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 5f7dc8b89511112aa467490b7874253665c4546e 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 +- clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp | 157 ++++++++++++------ .../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 9 +- clang/test/CIR/CodeGenOpenMP/parallel.c | 18 +- .../test/CIR/CodeGenOpenMP/target-parallel.c | 11 +- 5 files changed, 138 insertions(+), 72 deletions(-) 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/CIRGenStmtOpenMP.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp index 76e5be3614a1d..5075fdbfd5a7c 100644 --- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp @@ -47,26 +47,33 @@ getLeafClauses(CIRGenFunction &cgf, const OMPExecutableDirective &s, return result; } -template <typename DirectiveTy> +/// Evaluate the clauses allowed on an omp.parallel leaf into \p clauseOps. +/// Returns failure if a not-yet-implemented clause was encountered. 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); +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); +} + +/// Create an omp.parallel op from \p clauseOps and emit its body via +/// \p emitBody. +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 +83,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 +102,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 +304,18 @@ emitOMPTargetImplicitCaptures(CIRGenFunction &cgf, } } -template <typename DirectiveTy> +/// Evaluate the clauses allowed on an omp.target leaf into \p clauseOps, +/// collecting the mapped VarDecls into \p mapSyms. Returns failure if a +/// not-yet-implemented clause was encountered. 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 +323,17 @@ emitTargetOp(CIRGenFunction &cgf, const DirectiveTy &s, mlir::Location begin, OMPIsDevicePtrClause, OMPNowaitClause, OMPPrivateClause, OMPThreadLimitClause, OMPUsesAllocatorsClause, OMPXBareClause>{}, llvm::omp::Directive::OMPD_target); +} + +/// Create an omp.target op from \p clauseOps and emit its body via \p emitBody. +/// \p isCombined marks the op as a non-innermost leaf of a combined construct. +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, bool isCombined, + llvm::function_ref<mlir::LogicalResult()> emitBody) { + CIRGenBuilderTy &builder = cgf.getBuilder(); emitOMPTargetImplicitCaptures(cgf, s, mapSyms); @@ -317,6 +342,10 @@ emitTargetOp(CIRGenFunction &cgf, const DirectiveTy &s, mlir::Location begin, &cgf.getMLIRContext(), mlir::omp::TargetExecMode::generic); auto targetOp = mlir::omp::TargetOp::create(builder, begin, clauseOps); + // Mark the target as a non-innermost leaf of a combined construct so that + // later passes treat the nested leaves as a single combined construct. + if (isCombined) + targetOp.setCombined(true); mlir::Block &block = targetOp.getRegion().emplaceBlock(); for (mlir::Value mapVar : clauseOps.mapVars) @@ -349,10 +378,21 @@ 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, + /*isCombined=*/false, [&]() -> 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 +433,33 @@ 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, /*isCombined=*/true, + [&]() -> 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..bddbe62722f8d 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -5865,8 +5865,15 @@ 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()); + // Convert stack allocations used inside an omp.parallel nested in a generic + // omp.target region to device shared memory. This must run after the LLVM + // dialect lowering so all allocas are visible. The pass is anchored on + // LLVM::LLVMFuncOp, so nest it explicitly under the module pass manager. + pm.nest<mlir::LLVM::LLVMFuncOp>().addPass( + mlir::omp::createStackToSharedPass()); + } } std::unique_ptr<llvm::Module> diff --git a/clang/test/CIR/CodeGenOpenMP/parallel.c b/clang/test/CIR/CodeGenOpenMP/parallel.c index 36a48b38ee789..fd5dd0f1c30b1 100644 --- a/clang/test/CIR/CodeGenOpenMP/parallel.c +++ b/clang/test/CIR/CodeGenOpenMP/parallel.c @@ -36,25 +36,15 @@ 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. + // TODO(OMP): The shared and firstprivate clauses are not yet implemented, so + // codegen emits NYI diagnostics and skips the omp.parallel op entirely. When + // those clauses are implemented, the check lines will need updating. #pragma omp parallel shared(a) firstprivate(b) { a = a + 1; b = b + 1; } - // CHECK-NEXT: omp.parallel { - // CHECK-NEXT: cir.load align(4) %{{.*}} - // CHECK-NEXT: cir.const #cir.int<1> : !s32i - // CHECK-NEXT: cir.add nsw %{{.*}}, %{{.*}} : !s32i - // CHECK-NEXT: cir.store align(4) %{{.*}}, %{{.*}} : !s32i, !cir.ptr<!s32i> - // CHECK-NEXT: cir.load align(4) %{{.*}} - // CHECK-NEXT: cir.const #cir.int<1> : !s32i - // CHECK-NEXT: cir.add nsw %{{.*}}, %{{.*}} : !s32i - // CHECK-NEXT: cir.store align(4) %{{.*}}, %{{.*}} : !s32i, !cir.ptr<!s32i> - // CHECK-NEXT: omp.terminator - // CHECK-NEXT: } + // CHECK-NEXT: cir.return } void proc_bind_parallel() { // CHECK: cir.func{{.*}}@proc_bind_parallel 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
