https://github.com/jsjodin updated https://github.com/llvm/llvm-project/pull/195452
>From 46d10180ba08f3464a8ac75fa54cc791b577cc0e Mon Sep 17 00:00:00 2001 From: Jan Leyonberg <[email protected]> Date: Mon, 16 Mar 2026 12:07:44 -0400 Subject: [PATCH 1/6] [CIR][OpenMP] Initial implementation of target region support This patch adds support for target regions with some basic support for map clauses. It also changes the clause handling to make use of the OMP dialect ClauseOps to simplify op constrution. Assisted-by: Cursor / claude-4.6-opus-high --- clang/lib/CIR/CodeGen/CIRGenFunction.h | 4 - clang/lib/CIR/CodeGen/CIRGenOpenMPClause.cpp | 207 ++++++++++++------ clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h | 76 +++++++ clang/lib/CIR/CodeGen/CIRGenStmt.cpp | 3 +- clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp | 114 +++++++++- .../CIR/CodeGenOpenMP/not-yet-implemented.c | 4 +- .../CodeGenOpenMP/target-map-llvm-device.c | 111 ++++++++++ .../CIR/CodeGenOpenMP/target-map-llvm-host.c | 122 +++++++++++ clang/test/CIR/CodeGenOpenMP/target-map.c | 105 +++++++++ 9 files changed, 660 insertions(+), 86 deletions(-) create mode 100644 clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h create mode 100644 clang/test/CIR/CodeGenOpenMP/target-map-llvm-device.c create mode 100644 clang/test/CIR/CodeGenOpenMP/target-map-llvm-host.c create mode 100644 clang/test/CIR/CodeGenOpenMP/target-map.c diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h index 317151c8d61c6..88761fe91bc5e 100644 --- a/clang/lib/CIR/CodeGen/CIRGenFunction.h +++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h @@ -2559,10 +2559,6 @@ class CIRGenFunction : public CIRGenTypeCache { void emitOMPDeclareMapper(const OMPDeclareMapperDecl &d); void emitOMPRequiresDecl(const OMPRequiresDecl &d); -private: - template <typename Op> - void emitOpenMPClauses(Op &op, ArrayRef<const OMPClause *> clauses); - //===--------------------------------------------------------------------===// // OpenACC Emission //===--------------------------------------------------------------------===// diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.cpp index a0f0ea9299c8d..96194ad2316f5 100644 --- a/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.cpp @@ -6,90 +6,155 @@ // //===----------------------------------------------------------------------===// // -// Emit OpenMP clause nodes as CIR code. +// OpenMP clause processor implementation. See CIRGenOpenMPClause.h. // //===----------------------------------------------------------------------===// +#include "CIRGenOpenMPClause.h" #include "CIRGenFunction.h" #include "mlir/Dialect/OpenMP/OpenMPDialect.h" +#include "clang/Basic/OpenMPKinds.h" using namespace clang; using namespace clang::CIRGen; -namespace { -template <typename OpTy> -class OpenMPClauseCIREmitter final - : public ConstOMPClauseVisitor<OpenMPClauseCIREmitter<OpTy>> { - OpTy &operation; - CIRGen::CIRGenFunction &cgf; - CIRGen::CIRGenBuilderTy &builder; - -public: - OpenMPClauseCIREmitter(OpTy &operation, CIRGen::CIRGenFunction &cgf, - CIRGen::CIRGenBuilderTy &builder) - : operation(operation), cgf(cgf), builder(builder) {} - - void VisitOMPClause(const OMPClause *clause) { - cgf.cgm.errorNYI(clause->getBeginLoc(), "OpenMPClause ", - llvm::omp::getOpenMPClauseName(clause->getClauseKind())); +static mlir::omp::ClauseMapFlags +mapClauseKindToFlags(OpenMPMapClauseKind kind) { + switch (kind) { + case OMPC_MAP_to: + return mlir::omp::ClauseMapFlags::to; + case OMPC_MAP_from: + return mlir::omp::ClauseMapFlags::from; + case OMPC_MAP_tofrom: + return mlir::omp::ClauseMapFlags::to | mlir::omp::ClauseMapFlags::from; + case OMPC_MAP_alloc: + case OMPC_MAP_release: + return mlir::omp::ClauseMapFlags::storage; + case OMPC_MAP_delete: + return mlir::omp::ClauseMapFlags::del; + default: + return mlir::omp::ClauseMapFlags::none; } +} - void VisitOMPProcBindClause(const OMPProcBindClause *clause) { - if constexpr (std::is_same_v<OpTy, mlir::omp::ParallelOp>) { - mlir::omp::ClauseProcBindKind kind; - switch (clause->getProcBindKind()) { - case llvm::omp::ProcBindKind::OMP_PROC_BIND_master: - kind = mlir::omp::ClauseProcBindKind::Master; - break; - case llvm::omp::ProcBindKind::OMP_PROC_BIND_close: - kind = mlir::omp::ClauseProcBindKind::Close; - break; - case llvm::omp::ProcBindKind::OMP_PROC_BIND_spread: - kind = mlir::omp::ClauseProcBindKind::Spread; - break; - case llvm::omp::ProcBindKind::OMP_PROC_BIND_primary: - kind = mlir::omp::ClauseProcBindKind::Primary; - break; - case llvm::omp::ProcBindKind::OMP_PROC_BIND_default: - // 'default' in the classic-codegen does no runtime call/doesn't - // really do anything. So this is a no-op, and thus shouldn't change - // the IR. - return; - case llvm::omp::ProcBindKind::OMP_PROC_BIND_unknown: - llvm_unreachable("unknown proc-bind kind"); - } - operation.setProcBindKind(kind); - } else { - cgf.cgm.errorNYI( - clause->getBeginLoc(), - "OMPProcBindClause unimplemented on this directive kind"); - } - } +static mlir::Value emitMapInfoForVar(CIRGenFunction &cgf, + mlir::OpBuilder &builder, + mlir::Location loc, const VarDecl *vd, + mlir::omp::ClauseMapFlags mapFlags) { + Address addr = cgf.getAddrOfLocalVar(vd); + mlir::Value varPtr = addr.getPointer(); + auto varPtrType = mlir::cast<cir::PointerType>(varPtr.getType()); + mlir::Type elementType = varPtrType.getPointee(); - void emitClauses(ArrayRef<const OMPClause *> clauses) { - for (const auto *c : clauses) - this->Visit(c); + // Cast to generic pointer if needed. + if (varPtrType.getAddrSpace()) { + auto genericPtrType = + cir::PointerType::get(builder.getContext(), elementType); + varPtr = cir::CastOp::create(builder, loc, genericPtrType, + cir::CastKind::address_space, varPtr); + varPtrType = genericPtrType; } -}; -template <typename OpTy> -auto makeClauseEmitter(OpTy &op, CIRGen::CIRGenFunction &cgf, - CIRGen::CIRGenBuilderTy &builder) { - return OpenMPClauseCIREmitter<OpTy>(op, cgf, builder); + + return mlir::omp::MapInfoOp::create( + builder, loc, + /*omp_ptr=*/varPtrType, + /*var_ptr=*/varPtr, + /*var_ptr_type=*/mlir::TypeAttr::get(elementType), + /*map_type=*/builder.getAttr<mlir::omp::ClauseMapFlagsAttr>(mapFlags), + /*map_capture_type=*/ + builder.getAttr<mlir::omp::VariableCaptureKindAttr>( + mlir::omp::VariableCaptureKind::ByRef), + /*var_ptr_ptr=*/mlir::Value{}, + /*var_ptr_ptr_type=*/mlir::TypeAttr{}, + /*members=*/mlir::ValueRange{}, + /*members_index=*/mlir::ArrayAttr{}, + /*bounds=*/mlir::ValueRange{}, + /*mapper_id=*/mlir::FlatSymbolRefAttr{}, + /*name=*/builder.getStringAttr(vd->getName()), + /*partial_map=*/builder.getBoolAttr(false)); } -} // namespace - -template <typename Op> -void CIRGenFunction::emitOpenMPClauses(Op &op, - ArrayRef<const OMPClause *> clauses) { - mlir::OpBuilder::InsertionGuard guardCase(builder); - builder.setInsertionPoint(op); - makeClauseEmitter(op, *this, builder).emitClauses(clauses); + +bool OpenMPClauseProcessor::processProcBind( + mlir::omp::ProcBindClauseOps &result) const { + for (const OMPClause *clause : clauses) { + const auto *pbc = dyn_cast<OMPProcBindClause>(clause); + if (!pbc) + continue; + + switch (pbc->getProcBindKind()) { + case llvm::omp::ProcBindKind::OMP_PROC_BIND_master: + result.procBindKind = mlir::omp::ClauseProcBindKindAttr::get( + builder.getContext(), mlir::omp::ClauseProcBindKind::Master); + break; + case llvm::omp::ProcBindKind::OMP_PROC_BIND_close: + result.procBindKind = mlir::omp::ClauseProcBindKindAttr::get( + builder.getContext(), mlir::omp::ClauseProcBindKind::Close); + break; + case llvm::omp::ProcBindKind::OMP_PROC_BIND_spread: + result.procBindKind = mlir::omp::ClauseProcBindKindAttr::get( + builder.getContext(), mlir::omp::ClauseProcBindKind::Spread); + break; + case llvm::omp::ProcBindKind::OMP_PROC_BIND_primary: + result.procBindKind = mlir::omp::ClauseProcBindKindAttr::get( + builder.getContext(), mlir::omp::ClauseProcBindKind::Primary); + break; + case llvm::omp::ProcBindKind::OMP_PROC_BIND_default: + break; + case llvm::omp::ProcBindKind::OMP_PROC_BIND_unknown: + llvm_unreachable("unknown proc-bind kind"); + } + return true; + } + return false; } -// We're defining the template for this in a .cpp file, so we have to explicitly -// specialize the templates. -#define EXPL_SPEC(N) \ - template void CIRGenFunction::emitOpenMPClauses<N>( \ - N &, ArrayRef<const OMPClause *>); -EXPL_SPEC(mlir::omp::ParallelOp) -#undef EXPL_SPEC +bool OpenMPClauseProcessor::processMap( + mlir::omp::MapClauseOps &result, + llvm::SmallVectorImpl<const VarDecl *> *mapSyms) const { + bool found = false; + for (const OMPClause *clause : clauses) { + const auto *mc = dyn_cast<OMPMapClause>(clause); + if (!mc) + continue; + + found = true; + + for (OpenMPMapModifierKind mod : mc->getMapTypeModifiers()) { + if (mod != OMPC_MAP_MODIFIER_unknown) + cgm.errorNYI(mc->getBeginLoc(), + std::string("OpenMP map modifier '") + + getOpenMPSimpleClauseTypeName( + llvm::omp::Clause::OMPC_map, mod) + + "'"); + } + + if (mc->isImplicit()) { + cgm.errorNYI(mc->getBeginLoc(), "OpenMP implicit map clause"); + continue; + } + + mlir::omp::ClauseMapFlags mapFlags = mapClauseKindToFlags(mc->getMapType()); + + for (const Expr *varExpr : mc->varlist()) { + const auto *refExpr = dyn_cast<DeclRefExpr>(varExpr->IgnoreImplicit()); + if (!refExpr) { + cgm.errorNYI(varExpr->getExprLoc(), + "OpenMP map clause with non-DeclRefExpr variable"); + continue; + } + + const auto *vd = dyn_cast<VarDecl>(refExpr->getDecl()); + if (!vd) { + cgm.errorNYI(varExpr->getExprLoc(), + "OpenMP map clause with non-VarDecl variable"); + continue; + } + + result.mapVars.push_back( + emitMapInfoForVar(cgf, builder, loc, vd, mapFlags)); + if (mapSyms) + mapSyms->push_back(vd); + } + } + return found; +} diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h b/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h new file mode 100644 index 0000000000000..79f56c5bf12e9 --- /dev/null +++ b/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h @@ -0,0 +1,76 @@ +//===--- CIRGenOpenMPClause.h - OpenMP clause processor ---------*- 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_CIRGENOPENMPCLAUSE_H +#define LLVM_CLANG_LIB_CIR_CODEGEN_CIRGENOPENMPCLAUSE_H + +#include "CIRGenBuilder.h" +#include "CIRGenModule.h" +#include "mlir/Dialect/OpenMP/OpenMPClauseOperands.h" +#include "clang/AST/OpenMPClause.h" +#include "clang/AST/StmtOpenMP.h" +#include "llvm/Frontend/OpenMP/OMPConstants.h" + +namespace clang::CIRGen { + +class CIRGenFunction; + +/// Processes OpenMP clauses for a directive, writing results into the +/// auto-generated ClauseOps from the OMP dialect. +class OpenMPClauseProcessor { + CIRGenFunction &cgf; + CIRGenModule &cgm; + CIRGenBuilderTy &builder; + mlir::Location loc; + llvm::ArrayRef<const OMPClause *> clauses; + +public: + OpenMPClauseProcessor(CIRGenFunction &cgf, CIRGenModule &cgm, + CIRGenBuilderTy &builder, mlir::Location loc, + llvm::ArrayRef<const OMPClause *> clauses) + : cgf(cgf), cgm(cgm), builder(builder), loc(loc), clauses(clauses) {} + + bool processProcBind(mlir::omp::ProcBindClauseOps &result) const; + + /// Process map clauses. The optional \p mapSyms parameter collects the + /// VarDecls corresponding to each map operand. + bool + processMap(mlir::omp::MapClauseOps &result, + llvm::SmallVectorImpl<const VarDecl *> *mapSyms = nullptr) const; + + /// Emit an errorNYI for each clause of the given types if present. + template <typename... ClauseTypes> + void processTODO(llvm::omp::Directive directive) const; + +private: + template <typename ClauseType> + void processTODOClause(llvm::omp::Directive directive) const; +}; + +template <typename ClauseType> +void OpenMPClauseProcessor::processTODOClause( + llvm::omp::Directive directive) const { + for (const OMPClause *c : clauses) { + if (isa<ClauseType>(c)) { + std::string msg = + ("OpenMP " + llvm::omp::getOpenMPDirectiveName(directive) + " " + + llvm::omp::getOpenMPClauseName(c->getClauseKind()) + " clause") + .str(); + cgm.errorNYI(c->getBeginLoc(), msg); + } + } +} + +template <typename... ClauseTypes> +void OpenMPClauseProcessor::processTODO(llvm::omp::Directive directive) const { + (processTODOClause<ClauseTypes>(directive), ...); +} + +} // namespace clang::CIRGen + +#endif // LLVM_CLANG_LIB_CIR_CODEGEN_CIRGENOPENMPCLAUSE_H diff --git a/clang/lib/CIR/CodeGen/CIRGenStmt.cpp b/clang/lib/CIR/CodeGen/CIRGenStmt.cpp index 922140a93aa5a..2a9791329397b 100644 --- a/clang/lib/CIR/CodeGen/CIRGenStmt.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenStmt.cpp @@ -421,7 +421,6 @@ mlir::LogicalResult CIRGenFunction::emitStmt(const Stmt *s, case Stmt::CaseStmtClass: case Stmt::SEHLeaveStmtClass: case Stmt::SYCLKernelCallStmtClass: - case Stmt::CapturedStmtClass: case Stmt::ObjCAtTryStmtClass: case Stmt::ObjCAtThrowStmtClass: case Stmt::ObjCAtSynchronizedStmtClass: @@ -434,6 +433,8 @@ mlir::LogicalResult CIRGenFunction::emitStmt(const Stmt *s, cgm.errorNYI(s->getSourceRange(), std::string("emitStmt: ") + s->getStmtClassName()); return mlir::failure(); + case Stmt::CapturedStmtClass: + llvm_unreachable("CapturedStmt should be handled by the parent directive"); } llvm_unreachable("Unexpected statement class"); diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp index eb4934644b519..3c3938a3159e8 100644 --- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp @@ -12,7 +12,9 @@ #include "CIRGenBuilder.h" #include "CIRGenFunction.h" +#include "CIRGenOpenMPClause.h" #include "mlir/Dialect/OpenMP/OpenMPDialect.h" +#include "clang/AST/OpenMPClause.h" #include "clang/AST/StmtOpenMP.h" #include "llvm/Frontend/OpenMP/OMPConstants.h" using namespace clang; @@ -31,14 +33,19 @@ CIRGenFunction::emitOMPErrorDirective(const OMPErrorDirective &s) { mlir::LogicalResult CIRGenFunction::emitOMPParallelDirective(const OMPParallelDirective &s) { mlir::LogicalResult res = mlir::success(); - llvm::SmallVector<mlir::Type> retTy; - llvm::SmallVector<mlir::Value> operands; mlir::Location begin = getLoc(s.getBeginLoc()); mlir::Location end = getLoc(s.getEndLoc()); - auto parallelOp = - mlir::omp::ParallelOp::create(builder, begin, retTy, operands); - emitOpenMPClauses(parallelOp, s.clauses()); + mlir::omp::ParallelOperands clauseOps; + OpenMPClauseProcessor cp(*this, getCIRGenModule(), builder, begin, + s.clauses()); + cp.processProcBind(clauseOps); + cp.processTODO<OMPAllocateClause, OMPCopyinClause, OMPDefaultClause, + OMPFirstprivateClause, OMPIfClause, OMPNumThreadsClause, + OMPPrivateClause, OMPReductionClause, OMPSharedClause>( + llvm::omp::Directive::OMPD_parallel); + + auto parallelOp = mlir::omp::ParallelOp::create(builder, begin, clauseOps); { mlir::Block &block = parallelOp.getRegion().emplaceBlock(); @@ -207,10 +214,103 @@ CIRGenFunction::emitOMPAtomicDirective(const OMPAtomicDirective &s) { getCIRGenModule().errorNYI(s.getSourceRange(), "OpenMP OMPAtomicDirective"); 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) { - getCIRGenModule().errorNYI(s.getSourceRange(), "OpenMP OMPTargetDirective"); - return mlir::failure(); + mlir::Location begin = getLoc(s.getBeginLoc()); + mlir::Location end = getLoc(s.getEndLoc()); + + mlir::omp::TargetOperands clauseOps; + llvm::SmallVector<const VarDecl *> mapSyms; + + OpenMPClauseProcessor cp(*this, getCIRGenModule(), builder, begin, + s.clauses()); + cp.processMap(clauseOps, &mapSyms); + cp.processTODO<OMPAllocateClause, OMPDefaultClause, OMPDefaultmapClause, + OMPDependClause, OMPDeviceClause, OMPFirstprivateClause, + OMPHasDeviceAddrClause, OMPIfClause, OMPInReductionClause, + OMPIsDevicePtrClause, OMPNowaitClause, OMPPrivateClause, + OMPThreadLimitClause, OMPUsesAllocatorsClause, OMPXBareClause>( + llvm::omp::Directive::OMPD_target); + + emitOMPTargetImplicitCaptures(*this, s, mapSyms); + + auto targetOp = mlir::omp::TargetOp::create(builder, begin, clauseOps); + + return emitOMPTargetBody(*this, s, targetOp, clauseOps.mapVars, mapSyms, + begin, end); } mlir::LogicalResult CIRGenFunction::emitOMPTeamsDirective(const OMPTeamsDirective &s) { diff --git a/clang/test/CIR/CodeGenOpenMP/not-yet-implemented.c b/clang/test/CIR/CodeGenOpenMP/not-yet-implemented.c index 6d59f45d6e5e4..3502f0f291b22 100644 --- a/clang/test/CIR/CodeGenOpenMP/not-yet-implemented.c +++ b/clang/test/CIR/CodeGenOpenMP/not-yet-implemented.c @@ -10,9 +10,7 @@ void do_things() { {} int i; - // TODO(OMP): We might consider overloading operator<< for OMPClauseKind in - // the future if we want to improve this. - // expected-error@+1{{ClangIR code gen Not Yet Implemented: OpenMPClause : if}} + // expected-error@+1{{ClangIR code gen Not Yet Implemented: OpenMP parallel if clause}} #pragma omp parallel if(i) {} } diff --git a/clang/test/CIR/CodeGenOpenMP/target-map-llvm-device.c b/clang/test/CIR/CodeGenOpenMP/target-map-llvm-device.c new file mode 100644 index 0000000000000..25f80f20122a3 --- /dev/null +++ b/clang/test/CIR/CodeGenOpenMP/target-map-llvm-device.c @@ -0,0 +1,111 @@ +// Two-step host-BC -> device pipeline that mirrors the offloading driver. +// +// Step 1: Host compilation to bitcode (provides offload entry info to device pass). +// RUN: %clang_cc1 -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa \ +// RUN: -fclangir -emit-llvm-bc %s -o %t-cir-host.bc +// RUN: %clang_cc1 -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa \ +// RUN: -emit-llvm-bc %s -o %t-ogcg-host.bc +// +// Step 2: Device compilation using host BC. +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fopenmp -fopenmp-is-target-device \ +// RUN: -fopenmp-host-ir-file-path %t-cir-host.bc \ +// RUN: -fclangir -emit-llvm %s -o - \ +// RUN: | FileCheck %s --check-prefix=LLVM +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fopenmp -fopenmp-is-target-device \ +// RUN: -fopenmp-host-ir-file-path %t-ogcg-host.bc \ +// RUN: -emit-llvm %s -o - \ +// RUN: | FileCheck %s --check-prefix=OGCG + +void use(int); + +void target_map_to(int x) { +#pragma omp target map(to : x) + { + use(x); + } +} + +void target_map_from(int x) { +#pragma omp target map(from : x) + { + x = 42; + } +} + +void target_map_tofrom(int x) { +#pragma omp target map(tofrom : x) + { + x = x + 1; + } +} + +void target_map_multiple(int a, int b) { +#pragma omp target map(to : a) map(from : b) + { + b = a; + } +} + +// LLVM-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_target_map_to_l +// LLVM: call i32 @__kmpc_target_init( +// LLVM: user_code.entry: +// LLVM: %[[V:.*]] = load i32, ptr %{{.*}}, align 4 +// LLVM: call void @use(i32 {{.*}} %[[V]]) +// LLVM: call void @__kmpc_target_deinit() +// LLVM: ret void + +// LLVM-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_target_map_from_l +// LLVM: call i32 @__kmpc_target_init( +// LLVM: user_code.entry: +// LLVM: store i32 42, ptr %{{.*}}, align 4 +// LLVM: call void @__kmpc_target_deinit() +// LLVM: ret void + +// LLVM-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_target_map_tofrom_l +// LLVM: call i32 @__kmpc_target_init( +// LLVM: user_code.entry: +// LLVM: %[[LD:.*]] = load i32, ptr %{{.*}}, align 4 +// LLVM: %[[ADD:.*]] = add nsw i32 %[[LD]], 1 +// LLVM: store i32 %[[ADD]], ptr %{{.*}}, align 4 +// LLVM: call void @__kmpc_target_deinit() +// LLVM: ret void + +// LLVM-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_target_map_multiple_l +// LLVM: call i32 @__kmpc_target_init( +// LLVM: user_code.entry: +// LLVM: %[[A:.*]] = load i32, ptr %{{.*}}, align 4 +// LLVM: store i32 %[[A]], ptr %{{.*}}, align 4 +// LLVM: call void @__kmpc_target_deinit() +// LLVM: ret void + +// OGCG-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_target_map_to_l +// OGCG: call i32 @__kmpc_target_init( +// OGCG: user_code.entry: +// OGCG: %[[V:.*]] = load i32, ptr %{{.*}}, align 4 +// OGCG: call void @use(i32 {{.*}} %[[V]]) +// OGCG: call void @__kmpc_target_deinit() +// OGCG: ret void + +// OGCG-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_target_map_from_l +// OGCG: call i32 @__kmpc_target_init( +// OGCG: user_code.entry: +// OGCG: store i32 42, ptr %{{.*}}, align 4 +// OGCG: call void @__kmpc_target_deinit() +// OGCG: ret void + +// OGCG-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_target_map_tofrom_l +// OGCG: call i32 @__kmpc_target_init( +// OGCG: user_code.entry: +// OGCG: %[[LD:.*]] = load i32, ptr %{{.*}}, align 4 +// OGCG: %[[ADD:.*]] = add nsw i32 %[[LD]], 1 +// OGCG: store i32 %[[ADD]], ptr %{{.*}}, align 4 +// OGCG: call void @__kmpc_target_deinit() +// OGCG: ret void + +// OGCG-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_target_map_multiple_l +// OGCG: call i32 @__kmpc_target_init( +// OGCG: user_code.entry: +// OGCG: %[[A:.*]] = load i32, ptr %{{.*}}, align 4 +// OGCG: store i32 %[[A]], ptr %{{.*}}, align 4 +// OGCG: call void @__kmpc_target_deinit() +// OGCG: ret void diff --git a/clang/test/CIR/CodeGenOpenMP/target-map-llvm-host.c b/clang/test/CIR/CodeGenOpenMP/target-map-llvm-host.c new file mode 100644 index 0000000000000..2b1c314cc8c0b --- /dev/null +++ b/clang/test/CIR/CodeGenOpenMP/target-map-llvm-host.c @@ -0,0 +1,122 @@ +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fopenmp \ +// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -fclangir -emit-llvm %s -o - \ +// RUN: | FileCheck %s --check-prefix=LLVM +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fopenmp \ +// RUN: -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - \ +// RUN: | FileCheck %s --check-prefix=OGCG + +void use(int); + +void target_map_to(int x) { +#pragma omp target map(to : x) + { + use(x); + } +} + +void target_map_from(int x) { +#pragma omp target map(from : x) + { + x = 42; + } +} + +void target_map_tofrom(int x) { +#pragma omp target map(tofrom : x) + { + x = x + 1; + } +} + +void target_map_multiple(int a, int b) { +#pragma omp target map(to : a) map(from : b) + { + b = a; + } +} + +// Host wrappers + +// LLVM-LABEL: define {{.*}} void @target_map_to(i32 +// LLVM: call i32 @__tgt_target_kernel( +// LLVM: omp_offload.failed: +// LLVM: call void @__omp_offloading_{{.*}}_target_map_to_l + +// LLVM-LABEL: define {{.*}} void @target_map_from(i32 +// LLVM: call i32 @__tgt_target_kernel( +// LLVM: omp_offload.failed: +// LLVM: call void @__omp_offloading_{{.*}}_target_map_from_l + +// LLVM-LABEL: define {{.*}} void @target_map_tofrom(i32 +// LLVM: call i32 @__tgt_target_kernel( +// LLVM: omp_offload.failed: +// LLVM: call void @__omp_offloading_{{.*}}_target_map_tofrom_l + +// LLVM-LABEL: define {{.*}} void @target_map_multiple(i32 +// LLVM: call i32 @__tgt_target_kernel( +// LLVM: omp_offload.failed: +// LLVM: call void @__omp_offloading_{{.*}}_target_map_multiple_l + +// Outlined target functions + +// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_map_to_l +// LLVM: %[[V:.*]] = load i32, ptr %{{.*}}, align 4 +// LLVM: call void @use(i32 {{.*}} %[[V]]) +// LLVM: ret void + +// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_map_from_l +// LLVM: store i32 42, ptr %{{.*}}, align 4 +// LLVM: ret void + +// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_map_tofrom_l +// LLVM: %[[LD:.*]] = load i32, ptr %{{.*}}, align 4 +// LLVM: %[[ADD:.*]] = add nsw i32 %[[LD]], 1 +// LLVM: store i32 %[[ADD]], ptr %{{.*}}, align 4 +// LLVM: ret void + +// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_map_multiple_l +// LLVM: %[[A:.*]] = load i32, ptr %{{.*}}, align 4 +// LLVM: store i32 %[[A]], ptr %{{.*}}, align 4 +// LLVM: ret void + +// OGCG interleaves host wrapper and outlined function per target region. + +// OGCG-LABEL: define {{.*}} void @target_map_to(i32 +// OGCG: call i32 @__tgt_target_kernel( +// OGCG: omp_offload.failed: +// OGCG: call void @__omp_offloading_{{.*}}_target_map_to_l + +// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_map_to_l +// OGCG: %[[V:.*]] = load i32, ptr %{{.*}}, align 4 +// OGCG: call void @use(i32 {{.*}} %[[V]]) +// OGCG: ret void + +// OGCG-LABEL: define {{.*}} void @target_map_from(i32 +// OGCG: call i32 @__tgt_target_kernel( +// OGCG: omp_offload.failed: +// OGCG: call void @__omp_offloading_{{.*}}_target_map_from_l + +// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_map_from_l +// OGCG: store i32 42, ptr %{{.*}}, align 4 +// OGCG: ret void + +// OGCG-LABEL: define {{.*}} void @target_map_tofrom(i32 +// OGCG: call i32 @__tgt_target_kernel( +// OGCG: omp_offload.failed: +// OGCG: call void @__omp_offloading_{{.*}}_target_map_tofrom_l + +// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_map_tofrom_l +// OGCG: %[[LD:.*]] = load i32, ptr %{{.*}}, align 4 +// OGCG: %[[ADD:.*]] = add nsw i32 %[[LD]], 1 +// OGCG: store i32 %[[ADD]], ptr %{{.*}}, align 4 +// OGCG: ret void + +// OGCG-LABEL: define {{.*}} void @target_map_multiple(i32 +// OGCG: call i32 @__tgt_target_kernel( +// OGCG: omp_offload.failed: +// OGCG: call void @__omp_offloading_{{.*}}_target_map_multiple_l + +// OGCG-LABEL: define internal void @__omp_offloading_{{.*}}_target_map_multiple_l +// OGCG: %[[A:.*]] = load i32, ptr %{{.*}}, align 4 +// OGCG: store i32 %[[A]], ptr %{{.*}}, align 4 +// OGCG: ret void diff --git a/clang/test/CIR/CodeGenOpenMP/target-map.c b/clang/test/CIR/CodeGenOpenMP/target-map.c new file mode 100644 index 0000000000000..22d0ee811b91a --- /dev/null +++ b/clang/test/CIR/CodeGenOpenMP/target-map.c @@ -0,0 +1,105 @@ +// Host compilation (x86 host, AMDGPU offload target): no address space on allocas. +// 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 in private address space, addrspacecast for map info. +// 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); + +void target_map_to(int x) { + // CIR-HOST: cir.func{{.*}}@target_map_to + // CIR-HOST: %[[X_ALLOCA:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["x", init] + // CIR-HOST: %[[MAP:.*]] = omp.map.info var_ptr(%[[X_ALLOCA]] : !cir.ptr<!s32i>, !s32i) map_clauses(to) capture(ByRef) -> !cir.ptr<!s32i> {name = "x"} + // CIR-HOST-NEXT: omp.target map_entries(%[[MAP]] -> %[[ARG:.*]] : !cir.ptr<!s32i>) { + // CIR-HOST-NEXT: %[[LOAD:.*]] = cir.load align(4) %[[ARG]] + // CIR-HOST-NEXT: cir.call @use(%[[LOAD]]) + // CIR-HOST-NEXT: omp.terminator + // CIR-HOST-NEXT: } + + // CIR-DEVICE: cir.func{{.*}}@target_map_to + // CIR-DEVICE: %[[X_ALLOCA:.*]] = cir.alloca !s32i, !cir.ptr<!s32i, target_address_space(5)>, ["x", init] + // CIR-DEVICE: %[[CAST:.*]] = cir.cast address_space %[[X_ALLOCA]] : !cir.ptr<!s32i, target_address_space(5)> -> !cir.ptr<!s32i> + // CIR-DEVICE: %[[MAP:.*]] = omp.map.info var_ptr(%[[CAST]] : !cir.ptr<!s32i>, !s32i) map_clauses(to) capture(ByRef) -> !cir.ptr<!s32i> {name = "x"} + // CIR-DEVICE-NEXT: omp.target map_entries(%[[MAP]] -> %[[ARG:.*]] : !cir.ptr<!s32i>) { + // CIR-DEVICE: omp.terminator + // CIR-DEVICE-NEXT: } +#pragma omp target map(to : x) + { + use(x); + } +} + +void target_map_from(int x) { + // CIR-HOST: cir.func{{.*}}@target_map_from + // CIR-HOST: %[[X_ALLOCA:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["x", init] + // CIR-HOST: %[[MAP:.*]] = omp.map.info var_ptr(%[[X_ALLOCA]] : !cir.ptr<!s32i>, !s32i) map_clauses(from) capture(ByRef) -> !cir.ptr<!s32i> {name = "x"} + // CIR-HOST-NEXT: omp.target map_entries(%[[MAP]] -> %[[ARG:.*]] : !cir.ptr<!s32i>) { + // CIR-HOST-NEXT: %[[C42:.*]] = cir.const #cir.int<42> : !s32i + // CIR-HOST-NEXT: cir.store align(4) %[[C42]], %[[ARG]] + // CIR-HOST-NEXT: omp.terminator + // CIR-HOST-NEXT: } + + // CIR-DEVICE: cir.func{{.*}}@target_map_from + // CIR-DEVICE: %[[X_ALLOCA:.*]] = cir.alloca !s32i, !cir.ptr<!s32i, target_address_space(5)>, ["x", init] + // CIR-DEVICE: %[[CAST:.*]] = cir.cast address_space %[[X_ALLOCA]] : !cir.ptr<!s32i, target_address_space(5)> -> !cir.ptr<!s32i> + // CIR-DEVICE: %[[MAP:.*]] = omp.map.info var_ptr(%[[CAST]] : !cir.ptr<!s32i>, !s32i) map_clauses(from) capture(ByRef) -> !cir.ptr<!s32i> {name = "x"} + // CIR-DEVICE-NEXT: omp.target map_entries(%[[MAP]] -> %[[ARG:.*]] : !cir.ptr<!s32i>) { + // CIR-DEVICE: omp.terminator + // CIR-DEVICE-NEXT: } +#pragma omp target map(from : x) + { + x = 42; + } +} + +void target_map_tofrom(int x) { + // CIR-HOST: cir.func{{.*}}@target_map_tofrom + // CIR-HOST: %[[X_ALLOCA:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["x", init] + // CIR-HOST: %[[MAP:.*]] = omp.map.info var_ptr(%[[X_ALLOCA]] : !cir.ptr<!s32i>, !s32i) map_clauses(tofrom) capture(ByRef) -> !cir.ptr<!s32i> {name = "x"} + // CIR-HOST-NEXT: omp.target map_entries(%[[MAP]] -> %[[ARG:.*]] : !cir.ptr<!s32i>) { + // CIR-HOST: omp.terminator + // CIR-HOST-NEXT: } + + // CIR-DEVICE: cir.func{{.*}}@target_map_tofrom + // CIR-DEVICE: %[[X_ALLOCA:.*]] = cir.alloca !s32i, !cir.ptr<!s32i, target_address_space(5)>, ["x", init] + // CIR-DEVICE: %[[CAST:.*]] = cir.cast address_space %[[X_ALLOCA]] : !cir.ptr<!s32i, target_address_space(5)> -> !cir.ptr<!s32i> + // CIR-DEVICE: %[[MAP:.*]] = omp.map.info var_ptr(%[[CAST]] : !cir.ptr<!s32i>, !s32i) map_clauses(tofrom) capture(ByRef) -> !cir.ptr<!s32i> {name = "x"} + // CIR-DEVICE-NEXT: omp.target map_entries(%[[MAP]] -> %[[ARG:.*]] : !cir.ptr<!s32i>) { + // CIR-DEVICE: omp.terminator + // CIR-DEVICE-NEXT: } +#pragma omp target map(tofrom : x) + { + x = x + 1; + } +} + +void target_map_multiple(int a, int b) { + // CIR-HOST: cir.func{{.*}}@target_map_multiple + // CIR-HOST-DAG: %[[A_ALLOCA:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["a", init] + // CIR-HOST-DAG: %[[B_ALLOCA:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["b", init] + // CIR-HOST: %[[MAP_A:.*]] = omp.map.info var_ptr(%[[A_ALLOCA]] : !cir.ptr<!s32i>, !s32i) map_clauses(to) capture(ByRef) -> !cir.ptr<!s32i> {name = "a"} + // CIR-HOST-NEXT: %[[MAP_B:.*]] = omp.map.info var_ptr(%[[B_ALLOCA]] : !cir.ptr<!s32i>, !s32i) map_clauses(from) capture(ByRef) -> !cir.ptr<!s32i> {name = "b"} + // CIR-HOST-NEXT: omp.target map_entries(%[[MAP_A]] -> %[[ARG_A:.*]], %[[MAP_B]] -> %[[ARG_B:.*]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) { + // CIR-HOST: omp.terminator + // CIR-HOST-NEXT: } + + // CIR-DEVICE: cir.func{{.*}}@target_map_multiple + // CIR-DEVICE-DAG: %[[A_ALLOCA:.*]] = cir.alloca !s32i, !cir.ptr<!s32i, target_address_space(5)>, ["a", init] + // CIR-DEVICE-DAG: %[[B_ALLOCA:.*]] = cir.alloca !s32i, !cir.ptr<!s32i, target_address_space(5)>, ["b", init] + // CIR-DEVICE: %[[CAST_A:.*]] = cir.cast address_space %[[A_ALLOCA]] : !cir.ptr<!s32i, target_address_space(5)> -> !cir.ptr<!s32i> + // CIR-DEVICE: %[[MAP_A:.*]] = omp.map.info var_ptr(%[[CAST_A]] : !cir.ptr<!s32i>, !s32i) map_clauses(to) capture(ByRef) -> !cir.ptr<!s32i> {name = "a"} + // CIR-DEVICE: %[[CAST_B:.*]] = cir.cast address_space %[[B_ALLOCA]] : !cir.ptr<!s32i, target_address_space(5)> -> !cir.ptr<!s32i> + // CIR-DEVICE: %[[MAP_B:.*]] = omp.map.info var_ptr(%[[CAST_B]] : !cir.ptr<!s32i>, !s32i) map_clauses(from) capture(ByRef) -> !cir.ptr<!s32i> {name = "b"} + // CIR-DEVICE: omp.target map_entries(%[[MAP_A]] -> %[[ARG_A:.*]], %[[MAP_B]] -> %[[ARG_B:.*]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) { + // CIR-DEVICE: omp.terminator + // CIR-DEVICE-NEXT: } +#pragma omp target map(to : a) map(from : b) + { + b = a; + } +} + +// TODO: Test implicit mapping. Currently NYI >From 3353516c4b16da6db3548820c876393884fdc151 Mon Sep 17 00:00:00 2001 From: Jan Leyonberg <[email protected]> Date: Mon, 8 Jun 2026 13:14:53 -0400 Subject: [PATCH 2/6] Rename and use better NYI mechanism. --- clang/lib/CIR/CodeGen/CIRGenOpenMPClause.cpp | 6 +- clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h | 62 ++++++++++++-------- clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp | 30 +++++----- 3 files changed, 56 insertions(+), 42 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.cpp index 96194ad2316f5..c12d2fa68233e 100644 --- a/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.cpp @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// // -// OpenMP clause processor implementation. See CIRGenOpenMPClause.h. +// OpenMP clause emitter implementation. See CIRGenOpenMPClause.h. // //===----------------------------------------------------------------------===// @@ -74,7 +74,7 @@ static mlir::Value emitMapInfoForVar(CIRGenFunction &cgf, /*partial_map=*/builder.getBoolAttr(false)); } -bool OpenMPClauseProcessor::processProcBind( +bool OpenMPClauseEmitter::emitProcBind( mlir::omp::ProcBindClauseOps &result) const { for (const OMPClause *clause : clauses) { const auto *pbc = dyn_cast<OMPProcBindClause>(clause); @@ -108,7 +108,7 @@ bool OpenMPClauseProcessor::processProcBind( return false; } -bool OpenMPClauseProcessor::processMap( +bool OpenMPClauseEmitter::emitMap( mlir::omp::MapClauseOps &result, llvm::SmallVectorImpl<const VarDecl *> *mapSyms) const { bool found = false; diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h b/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h index 79f56c5bf12e9..19649b265bc1b 100644 --- a/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h +++ b/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h @@ -1,4 +1,4 @@ -//===--- CIRGenOpenMPClause.h - OpenMP clause processor ---------*- C++ -*-===// +//===--- CIRGenOpenMPClause.h - OpenMP clause emitter -----------*- C++ -*-===// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. @@ -16,13 +16,20 @@ #include "clang/AST/StmtOpenMP.h" #include "llvm/Frontend/OpenMP/OMPConstants.h" +#include <type_traits> + namespace clang::CIRGen { class CIRGenFunction; -/// Processes OpenMP clauses for a directive, writing results into the +/// A type-only list of OpenMP clause AST node types. +/// Note: The clause AST classes do not have a default constructor, so a +/// std::tuple is not practical. +template <typename... Clauses> struct OpenMPClauseList {}; + +/// Emits OpenMP clauses for a directive, writing results into the /// auto-generated ClauseOps from the OMP dialect. -class OpenMPClauseProcessor { +class OpenMPClauseEmitter { CIRGenFunction &cgf; CIRGenModule &cgm; CIRGenBuilderTy &builder; @@ -30,47 +37,52 @@ class OpenMPClauseProcessor { llvm::ArrayRef<const OMPClause *> clauses; public: - OpenMPClauseProcessor(CIRGenFunction &cgf, CIRGenModule &cgm, - CIRGenBuilderTy &builder, mlir::Location loc, - llvm::ArrayRef<const OMPClause *> clauses) + OpenMPClauseEmitter(CIRGenFunction &cgf, CIRGenModule &cgm, + CIRGenBuilderTy &builder, mlir::Location loc, + llvm::ArrayRef<const OMPClause *> clauses) : cgf(cgf), cgm(cgm), builder(builder), loc(loc), clauses(clauses) {} - bool processProcBind(mlir::omp::ProcBindClauseOps &result) const; + bool emitProcBind(mlir::omp::ProcBindClauseOps &result) const; - /// Process map clauses. The optional \p mapSyms parameter collects the + /// Emit map clauses. The optional \p mapSyms parameter collects the /// VarDecls corresponding to each map operand. - bool - processMap(mlir::omp::MapClauseOps &result, - llvm::SmallVectorImpl<const VarDecl *> *mapSyms = nullptr) const; + bool emitMap(mlir::omp::MapClauseOps &result, + llvm::SmallVectorImpl<const VarDecl *> *mapSyms = nullptr) const; - /// Emit an errorNYI for each clause of the given types if present. - template <typename... ClauseTypes> - void processTODO(llvm::omp::Directive directive) const; + /// Verify the clauses of a directive to make sure all legal cases are either + /// implemented or give a NYI error. If the clause is neither, then + /// an unknown clause error will be emitted. + template <typename... SupportedClauses, typename... NYIClauses> + void emitNYI(OpenMPClauseList<NYIClauses...> nyi, + llvm::omp::Directive directive) const; private: - template <typename ClauseType> - void processTODOClause(llvm::omp::Directive directive) const; + /// True if T is the same type as any of Ts. + template <typename T, typename... Ts> + static constexpr bool isAnyOf = (std::is_same_v<T, Ts> || ...); }; -template <typename ClauseType> -void OpenMPClauseProcessor::processTODOClause( - llvm::omp::Directive directive) const { +template <typename... SupportedClauses, typename... NYIClauses> +void OpenMPClauseEmitter::emitNYI(OpenMPClauseList<NYIClauses...>, + llvm::omp::Directive directive) const { + static_assert( + (!isAnyOf<NYIClauses, SupportedClauses...> && ...), + "the supported and not-yet-implemented clause lists must be disjoint"); + for (const OMPClause *c : clauses) { - if (isa<ClauseType>(c)) { + if ((isa<NYIClauses>(c) || ...)) { std::string msg = ("OpenMP " + llvm::omp::getOpenMPDirectiveName(directive) + " " + llvm::omp::getOpenMPClauseName(c->getClauseKind()) + " clause") .str(); cgm.errorNYI(c->getBeginLoc(), msg); + } else if (!(isa<SupportedClauses>(c) || ...)) { + // Unknown/illegal clause encountered + llvm_unreachable("unexpected OpenMP clause"); } } } -template <typename... ClauseTypes> -void OpenMPClauseProcessor::processTODO(llvm::omp::Directive directive) const { - (processTODOClause<ClauseTypes>(directive), ...); -} - } // namespace clang::CIRGen #endif // LLVM_CLANG_LIB_CIR_CODEGEN_CIRGENOPENMPCLAUSE_H diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp index 3c3938a3159e8..38a55fbdb8c2a 100644 --- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp @@ -37,12 +37,13 @@ CIRGenFunction::emitOMPParallelDirective(const OMPParallelDirective &s) { mlir::Location end = getLoc(s.getEndLoc()); mlir::omp::ParallelOperands clauseOps; - OpenMPClauseProcessor cp(*this, getCIRGenModule(), builder, begin, - s.clauses()); - cp.processProcBind(clauseOps); - cp.processTODO<OMPAllocateClause, OMPCopyinClause, OMPDefaultClause, - OMPFirstprivateClause, OMPIfClause, OMPNumThreadsClause, - OMPPrivateClause, OMPReductionClause, OMPSharedClause>( + OpenMPClauseEmitter ce(*this, getCIRGenModule(), builder, begin, s.clauses()); + ce.emitProcBind(clauseOps); + ce.emitNYI</*supported=*/OMPProcBindClause>( + /*nyi=*/OpenMPClauseList< + OMPAllocateClause, OMPCopyinClause, OMPDefaultClause, + OMPFirstprivateClause, OMPIfClause, OMPNumThreadsClause, + OMPPrivateClause, OMPReductionClause, OMPSharedClause>{}, llvm::omp::Directive::OMPD_parallel); auto parallelOp = mlir::omp::ParallelOp::create(builder, begin, clauseOps); @@ -295,14 +296,15 @@ CIRGenFunction::emitOMPTargetDirective(const OMPTargetDirective &s) { mlir::omp::TargetOperands clauseOps; llvm::SmallVector<const VarDecl *> mapSyms; - OpenMPClauseProcessor cp(*this, getCIRGenModule(), builder, begin, - s.clauses()); - cp.processMap(clauseOps, &mapSyms); - cp.processTODO<OMPAllocateClause, OMPDefaultClause, OMPDefaultmapClause, - OMPDependClause, OMPDeviceClause, OMPFirstprivateClause, - OMPHasDeviceAddrClause, OMPIfClause, OMPInReductionClause, - OMPIsDevicePtrClause, OMPNowaitClause, OMPPrivateClause, - OMPThreadLimitClause, OMPUsesAllocatorsClause, OMPXBareClause>( + OpenMPClauseEmitter ce(*this, getCIRGenModule(), builder, begin, s.clauses()); + ce.emitMap(clauseOps, &mapSyms); + ce.emitNYI</*supported=*/OMPMapClause>( + /*nyi=*/OpenMPClauseList< + OMPAllocateClause, OMPDefaultClause, OMPDefaultmapClause, + OMPDependClause, OMPDeviceClause, OMPFirstprivateClause, + OMPHasDeviceAddrClause, OMPIfClause, OMPInReductionClause, + OMPIsDevicePtrClause, OMPNowaitClause, OMPPrivateClause, + OMPThreadLimitClause, OMPUsesAllocatorsClause, OMPXBareClause>{}, llvm::omp::Directive::OMPD_target); emitOMPTargetImplicitCaptures(*this, s, mapSyms); >From 3c648535272b70a91cad865cc6d267e9b8c43c92 Mon Sep 17 00:00:00 2001 From: Jan Leyonberg <[email protected]> Date: Mon, 8 Jun 2026 14:01:54 -0400 Subject: [PATCH 3/6] Create helper function --- clang/lib/CIR/CodeGen/CIRGenOpenMPClause.cpp | 27 +++++-------------- .../mlir/Dialect/OpenMP/OpenMPDialect.h | 6 +++++ mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp | 18 +++++++++++++ 3 files changed, 30 insertions(+), 21 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.cpp index c12d2fa68233e..84ab195c2aa52 100644 --- a/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.cpp @@ -81,28 +81,13 @@ bool OpenMPClauseEmitter::emitProcBind( if (!pbc) continue; - switch (pbc->getProcBindKind()) { - case llvm::omp::ProcBindKind::OMP_PROC_BIND_master: + llvm::omp::ProcBindKind kind = pbc->getProcBindKind(); + assert(kind != llvm::omp::ProcBindKind::OMP_PROC_BIND_unknown && + "unknown proc-bind kind"); + // The 'default' kind has no dialect counterpart; leave the attribute unset. + if (kind != llvm::omp::ProcBindKind::OMP_PROC_BIND_default) result.procBindKind = mlir::omp::ClauseProcBindKindAttr::get( - builder.getContext(), mlir::omp::ClauseProcBindKind::Master); - break; - case llvm::omp::ProcBindKind::OMP_PROC_BIND_close: - result.procBindKind = mlir::omp::ClauseProcBindKindAttr::get( - builder.getContext(), mlir::omp::ClauseProcBindKind::Close); - break; - case llvm::omp::ProcBindKind::OMP_PROC_BIND_spread: - result.procBindKind = mlir::omp::ClauseProcBindKindAttr::get( - builder.getContext(), mlir::omp::ClauseProcBindKind::Spread); - break; - case llvm::omp::ProcBindKind::OMP_PROC_BIND_primary: - result.procBindKind = mlir::omp::ClauseProcBindKindAttr::get( - builder.getContext(), mlir::omp::ClauseProcBindKind::Primary); - break; - case llvm::omp::ProcBindKind::OMP_PROC_BIND_default: - break; - case llvm::omp::ProcBindKind::OMP_PROC_BIND_unknown: - llvm_unreachable("unknown proc-bind kind"); - } + builder.getContext(), mlir::omp::convertProcBindKind(kind)); return true; } return false; diff --git a/mlir/include/mlir/Dialect/OpenMP/OpenMPDialect.h b/mlir/include/mlir/Dialect/OpenMP/OpenMPDialect.h index d7bccb133e02c..c61e9415327f0 100644 --- a/mlir/include/mlir/Dialect/OpenMP/OpenMPDialect.h +++ b/mlir/include/mlir/Dialect/OpenMP/OpenMPDialect.h @@ -24,6 +24,7 @@ #include "mlir/IR/SymbolTable.h" #include "mlir/Interfaces/ControlFlowInterfaces.h" #include "mlir/Interfaces/SideEffectInterfaces.h" +#include "llvm/Frontend/OpenMP/OMPConstants.h" #include "llvm/Frontend/OpenMP/OMPDeviceConstants.h" #define GET_TYPEDEF_CLASSES @@ -51,6 +52,11 @@ namespace mlir::omp { /// Find the omp.new_cli, generator, and consumer of a canonical loop info. std::tuple<NewCliOp, OpOperand *, OpOperand *> decodeCli(mlir::Value cli); + +/// Convert a proc_bind kind from the LLVM frontend enum to the corresponding +/// OpenMP dialect enum. The LLVM 'default' and 'unknown' kinds have no dialect +/// counterpart and are not valid inputs. +ClauseProcBindKind convertProcBindKind(llvm::omp::ProcBindKind kind); } // namespace mlir::omp #endif // MLIR_DIALECT_OPENMP_OPENMPDIALECT_H_ diff --git a/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp b/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp index c2956a1cf7b79..cbddb96713874 100644 --- a/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp +++ b/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp @@ -3945,6 +3945,24 @@ mlir::omp ::decodeCli(Value cli) { return {create, gen, cons}; } +ClauseProcBindKind +mlir::omp::convertProcBindKind(llvm::omp::ProcBindKind kind) { + switch (kind) { + case llvm::omp::ProcBindKind::OMP_PROC_BIND_close: + return ClauseProcBindKind::Close; + case llvm::omp::ProcBindKind::OMP_PROC_BIND_master: + return ClauseProcBindKind::Master; + case llvm::omp::ProcBindKind::OMP_PROC_BIND_primary: + return ClauseProcBindKind::Primary; + case llvm::omp::ProcBindKind::OMP_PROC_BIND_spread: + return ClauseProcBindKind::Spread; + case llvm::omp::ProcBindKind::OMP_PROC_BIND_default: + case llvm::omp::ProcBindKind::OMP_PROC_BIND_unknown: + break; + } + llvm_unreachable("unexpected proc-bind kind"); +} + void NewCliOp::build(::mlir::OpBuilder &odsBuilder, ::mlir::OperationState &odsState) { odsState.addTypes(CanonicalLoopInfoType::get(odsBuilder.getContext())); >From 3d31612e4221212f2cfd1c6b4d3547e5057fdf5b Mon Sep 17 00:00:00 2001 From: Jan Leyonberg <[email protected]> Date: Mon, 8 Jun 2026 14:32:19 -0400 Subject: [PATCH 4/6] Fix other review comments. --- clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h | 5 +- clang/lib/CIR/CodeGen/CIRGenStmt.cpp | 2 +- .../CIR/CodeGenOpenMP/not-yet-implemented.c | 2 +- .../CodeGenOpenMP/target-map-llvm-device.c | 31 ++++++-- .../CIR/CodeGenOpenMP/target-map-llvm-host.c | 77 ++++++++++++++----- 5 files changed, 89 insertions(+), 28 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h b/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h index 19649b265bc1b..c1fa31ff2bd11 100644 --- a/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h +++ b/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h @@ -72,8 +72,9 @@ void OpenMPClauseEmitter::emitNYI(OpenMPClauseList<NYIClauses...>, for (const OMPClause *c : clauses) { if ((isa<NYIClauses>(c) || ...)) { std::string msg = - ("OpenMP " + llvm::omp::getOpenMPDirectiveName(directive) + " " + - llvm::omp::getOpenMPClauseName(c->getClauseKind()) + " clause") + (llvm::Twine("OpenMP ") + + llvm::omp::getOpenMPDirectiveName(directive).upper() + " '" + + llvm::omp::getOpenMPClauseName(c->getClauseKind()) + "' clause") .str(); cgm.errorNYI(c->getBeginLoc(), msg); } else if (!(isa<SupportedClauses>(c) || ...)) { diff --git a/clang/lib/CIR/CodeGen/CIRGenStmt.cpp b/clang/lib/CIR/CodeGen/CIRGenStmt.cpp index 2a9791329397b..955163c23e32e 100644 --- a/clang/lib/CIR/CodeGen/CIRGenStmt.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenStmt.cpp @@ -434,7 +434,7 @@ mlir::LogicalResult CIRGenFunction::emitStmt(const Stmt *s, std::string("emitStmt: ") + s->getStmtClassName()); return mlir::failure(); case Stmt::CapturedStmtClass: - llvm_unreachable("CapturedStmt should be handled by the parent directive"); + llvm_unreachable("CapturedStmt must be handled by the parent directive"); } llvm_unreachable("Unexpected statement class"); diff --git a/clang/test/CIR/CodeGenOpenMP/not-yet-implemented.c b/clang/test/CIR/CodeGenOpenMP/not-yet-implemented.c index 3502f0f291b22..29e280bf262d0 100644 --- a/clang/test/CIR/CodeGenOpenMP/not-yet-implemented.c +++ b/clang/test/CIR/CodeGenOpenMP/not-yet-implemented.c @@ -10,7 +10,7 @@ void do_things() { {} int i; - // expected-error@+1{{ClangIR code gen Not Yet Implemented: OpenMP parallel if clause}} + // expected-error@+1{{ClangIR code gen Not Yet Implemented: OpenMP PARALLEL 'if' clause}} #pragma omp parallel if(i) {} } diff --git a/clang/test/CIR/CodeGenOpenMP/target-map-llvm-device.c b/clang/test/CIR/CodeGenOpenMP/target-map-llvm-device.c index 25f80f20122a3..b6fa078d143c5 100644 --- a/clang/test/CIR/CodeGenOpenMP/target-map-llvm-device.c +++ b/clang/test/CIR/CodeGenOpenMP/target-map-llvm-device.c @@ -47,34 +47,53 @@ void target_map_multiple(int a, int b) { } // LLVM-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_target_map_to_l +// LLVM-SAME: (ptr %[[ARG:[^,]+]], ptr +// LLVM: %[[SLOT:.*]] = addrspacecast ptr addrspace(5) %{{.*}} to ptr +// LLVM: store ptr %[[ARG]], ptr %[[SLOT]], align 8 // LLVM: call i32 @__kmpc_target_init( // LLVM: user_code.entry: -// LLVM: %[[V:.*]] = load i32, ptr %{{.*}}, align 4 +// LLVM: %[[PTR:.*]] = load ptr, ptr %[[SLOT]], align 8 +// LLVM: %[[V:.*]] = load i32, ptr %[[PTR]], align 4 // LLVM: call void @use(i32 {{.*}} %[[V]]) // LLVM: call void @__kmpc_target_deinit() // LLVM: ret void // LLVM-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_target_map_from_l +// LLVM-SAME: (ptr %[[ARG:[^,]+]], ptr +// LLVM: %[[SLOT:.*]] = addrspacecast ptr addrspace(5) %{{.*}} to ptr +// LLVM: store ptr %[[ARG]], ptr %[[SLOT]], align 8 // LLVM: call i32 @__kmpc_target_init( // LLVM: user_code.entry: -// LLVM: store i32 42, ptr %{{.*}}, align 4 +// LLVM: %[[PTR:.*]] = load ptr, ptr %[[SLOT]], align 8 +// LLVM: store i32 42, ptr %[[PTR]], align 4 // LLVM: call void @__kmpc_target_deinit() // LLVM: ret void // LLVM-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_target_map_tofrom_l +// LLVM-SAME: (ptr %[[ARG:[^,]+]], ptr +// LLVM: %[[SLOT:.*]] = addrspacecast ptr addrspace(5) %{{.*}} to ptr +// LLVM: store ptr %[[ARG]], ptr %[[SLOT]], align 8 // LLVM: call i32 @__kmpc_target_init( // LLVM: user_code.entry: -// LLVM: %[[LD:.*]] = load i32, ptr %{{.*}}, align 4 +// LLVM: %[[PTR:.*]] = load ptr, ptr %[[SLOT]], align 8 +// LLVM: %[[LD:.*]] = load i32, ptr %[[PTR]], align 4 // LLVM: %[[ADD:.*]] = add nsw i32 %[[LD]], 1 -// LLVM: store i32 %[[ADD]], ptr %{{.*}}, align 4 +// LLVM: store i32 %[[ADD]], ptr %[[PTR]], align 4 // LLVM: call void @__kmpc_target_deinit() // LLVM: ret void // LLVM-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_target_map_multiple_l +// LLVM-SAME: (ptr %[[ARG_A:[^,]+]], ptr %[[ARG_B:[^,]+]], ptr +// LLVM: %[[SLOT_A:.*]] = addrspacecast ptr addrspace(5) %{{.*}} to ptr +// LLVM: store ptr %[[ARG_A]], ptr %[[SLOT_A]], align 8 +// LLVM: %[[SLOT_B:.*]] = addrspacecast ptr addrspace(5) %{{.*}} to ptr +// LLVM: store ptr %[[ARG_B]], ptr %[[SLOT_B]], align 8 // LLVM: call i32 @__kmpc_target_init( // LLVM: user_code.entry: -// LLVM: %[[A:.*]] = load i32, ptr %{{.*}}, align 4 -// LLVM: store i32 %[[A]], ptr %{{.*}}, align 4 +// LLVM: %[[PTR_A:.*]] = load ptr, ptr %[[SLOT_A]], align 8 +// LLVM: %[[PTR_B:.*]] = load ptr, ptr %[[SLOT_B]], align 8 +// LLVM: %[[A:.*]] = load i32, ptr %[[PTR_A]], align 4 +// LLVM: store i32 %[[A]], ptr %[[PTR_B]], align 4 // LLVM: call void @__kmpc_target_deinit() // LLVM: ret void diff --git a/clang/test/CIR/CodeGenOpenMP/target-map-llvm-host.c b/clang/test/CIR/CodeGenOpenMP/target-map-llvm-host.c index 2b1c314cc8c0b..a42e955396980 100644 --- a/clang/test/CIR/CodeGenOpenMP/target-map-llvm-host.c +++ b/clang/test/CIR/CodeGenOpenMP/target-map-llvm-host.c @@ -36,47 +36,88 @@ void target_map_multiple(int a, int b) { } // Host wrappers - -// LLVM-LABEL: define {{.*}} void @target_map_to(i32 +// +// LLVM-LABEL: define {{.*}} void @target_map_to( +// LLVM-SAME: i32 noundef %[[ARG:[^,)]+]] +// LLVM: %[[X_ADDR:.*]] = alloca i32, i64 1, align 4 +// LLVM: store i32 %[[ARG]], ptr %[[X_ADDR]], align 4 +// LLVM: %[[BP:.*]] = getelementptr inbounds [2 x ptr], ptr %.offload_baseptrs, i32 0, i32 0 +// LLVM: store ptr %[[X_ADDR]], ptr %[[BP]], align 8 +// LLVM: %[[P:.*]] = getelementptr inbounds [2 x ptr], ptr %.offload_ptrs, i32 0, i32 0 +// LLVM: store ptr %[[X_ADDR]], ptr %[[P]], align 8 // LLVM: call i32 @__tgt_target_kernel( // LLVM: omp_offload.failed: -// LLVM: call void @__omp_offloading_{{.*}}_target_map_to_l - -// LLVM-LABEL: define {{.*}} void @target_map_from(i32 +// LLVM: call void @__omp_offloading_{{.*}}_target_map_to_l{{.*}}(ptr %[[X_ADDR]], ptr null) + +// LLVM-LABEL: define {{.*}} void @target_map_from( +// LLVM-SAME: i32 noundef %[[ARG:[^,)]+]] +// LLVM: %[[X_ADDR:.*]] = alloca i32, i64 1, align 4 +// LLVM: store i32 %[[ARG]], ptr %[[X_ADDR]], align 4 +// LLVM: %[[BP:.*]] = getelementptr inbounds [2 x ptr], ptr %.offload_baseptrs, i32 0, i32 0 +// LLVM: store ptr %[[X_ADDR]], ptr %[[BP]], align 8 +// LLVM: %[[P:.*]] = getelementptr inbounds [2 x ptr], ptr %.offload_ptrs, i32 0, i32 0 +// LLVM: store ptr %[[X_ADDR]], ptr %[[P]], align 8 // LLVM: call i32 @__tgt_target_kernel( // LLVM: omp_offload.failed: -// LLVM: call void @__omp_offloading_{{.*}}_target_map_from_l - -// LLVM-LABEL: define {{.*}} void @target_map_tofrom(i32 +// LLVM: call void @__omp_offloading_{{.*}}_target_map_from_l{{.*}}(ptr %[[X_ADDR]], ptr null) + +// LLVM-LABEL: define {{.*}} void @target_map_tofrom( +// LLVM-SAME: i32 noundef %[[ARG:[^,)]+]] +// LLVM: %[[X_ADDR:.*]] = alloca i32, i64 1, align 4 +// LLVM: store i32 %[[ARG]], ptr %[[X_ADDR]], align 4 +// LLVM: %[[BP:.*]] = getelementptr inbounds [2 x ptr], ptr %.offload_baseptrs, i32 0, i32 0 +// LLVM: store ptr %[[X_ADDR]], ptr %[[BP]], align 8 +// LLVM: %[[P:.*]] = getelementptr inbounds [2 x ptr], ptr %.offload_ptrs, i32 0, i32 0 +// LLVM: store ptr %[[X_ADDR]], ptr %[[P]], align 8 // LLVM: call i32 @__tgt_target_kernel( // LLVM: omp_offload.failed: -// LLVM: call void @__omp_offloading_{{.*}}_target_map_tofrom_l - -// LLVM-LABEL: define {{.*}} void @target_map_multiple(i32 +// LLVM: call void @__omp_offloading_{{.*}}_target_map_tofrom_l{{.*}}(ptr %[[X_ADDR]], ptr null) + +// LLVM-LABEL: define {{.*}} void @target_map_multiple( +// LLVM-SAME: i32 noundef %[[ARG_A:[^,)]+]], i32 noundef %[[ARG_B:[^,)]+]] +// LLVM: %[[A_ADDR:.*]] = alloca i32, i64 1, align 4 +// LLVM: %[[B_ADDR:.*]] = alloca i32, i64 1, align 4 +// LLVM: store i32 %[[ARG_A]], ptr %[[A_ADDR]], align 4 +// LLVM: store i32 %[[ARG_B]], ptr %[[B_ADDR]], align 4 +// LLVM: %[[BP_A:.*]] = getelementptr inbounds [3 x ptr], ptr %.offload_baseptrs, i32 0, i32 0 +// LLVM: store ptr %[[A_ADDR]], ptr %[[BP_A]], align 8 +// LLVM: %[[P_A:.*]] = getelementptr inbounds [3 x ptr], ptr %.offload_ptrs, i32 0, i32 0 +// LLVM: store ptr %[[A_ADDR]], ptr %[[P_A]], align 8 +// LLVM: %[[BP_B:.*]] = getelementptr inbounds [3 x ptr], ptr %.offload_baseptrs, i32 0, i32 1 +// LLVM: store ptr %[[B_ADDR]], ptr %[[BP_B]], align 8 +// LLVM: %[[P_B:.*]] = getelementptr inbounds [3 x ptr], ptr %.offload_ptrs, i32 0, i32 1 +// LLVM: store ptr %[[B_ADDR]], ptr %[[P_B]], align 8 // LLVM: call i32 @__tgt_target_kernel( // LLVM: omp_offload.failed: -// LLVM: call void @__omp_offloading_{{.*}}_target_map_multiple_l +// LLVM: call void @__omp_offloading_{{.*}}_target_map_multiple_l{{.*}}(ptr %[[A_ADDR]], ptr %[[B_ADDR]], ptr null) // Outlined target functions +// +// The mapped pointer arrives as the first function argument; load/store the +// user value directly through it. // LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_map_to_l -// LLVM: %[[V:.*]] = load i32, ptr %{{.*}}, align 4 +// LLVM-SAME: (ptr %[[ARG:[^,]+]], ptr +// LLVM: %[[V:.*]] = load i32, ptr %[[ARG]], align 4 // LLVM: call void @use(i32 {{.*}} %[[V]]) // LLVM: ret void // LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_map_from_l -// LLVM: store i32 42, ptr %{{.*}}, align 4 +// LLVM-SAME: (ptr %[[ARG:[^,]+]], ptr +// LLVM: store i32 42, ptr %[[ARG]], align 4 // LLVM: ret void // LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_map_tofrom_l -// LLVM: %[[LD:.*]] = load i32, ptr %{{.*}}, align 4 +// LLVM-SAME: (ptr %[[ARG:[^,]+]], ptr +// LLVM: %[[LD:.*]] = load i32, ptr %[[ARG]], align 4 // LLVM: %[[ADD:.*]] = add nsw i32 %[[LD]], 1 -// LLVM: store i32 %[[ADD]], ptr %{{.*}}, align 4 +// LLVM: store i32 %[[ADD]], ptr %[[ARG]], align 4 // LLVM: ret void // LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_map_multiple_l -// LLVM: %[[A:.*]] = load i32, ptr %{{.*}}, align 4 -// LLVM: store i32 %[[A]], ptr %{{.*}}, align 4 +// LLVM-SAME: (ptr %[[ARG_A:[^,]+]], ptr %[[ARG_B:[^,]+]], ptr +// LLVM: %[[A:.*]] = load i32, ptr %[[ARG_A]], align 4 +// LLVM: store i32 %[[A]], ptr %[[ARG_B]], align 4 // LLVM: ret void // OGCG interleaves host wrapper and outlined function per target region. >From 34ce35b1236001e48b5294b50ca5274ab02e553b Mon Sep 17 00:00:00 2001 From: Jan Leyonberg <[email protected]> Date: Mon, 15 Jun 2026 13:16:47 -0400 Subject: [PATCH 5/6] Review feedback fixes. --- clang/lib/CIR/CodeGen/CIRGenOpenMPClause.cpp | 2 +- clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h | 15 +++++++-------- clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp | 4 ++-- 3 files changed, 10 insertions(+), 11 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.cpp index 84ab195c2aa52..16ac4440660b5 100644 --- a/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.cpp @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// // -// OpenMP clause emitter implementation. See CIRGenOpenMPClause.h. +// OpenMP clause emitter implementation. // //===----------------------------------------------------------------------===// diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h b/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h index c1fa31ff2bd11..a28fd6439bfdd 100644 --- a/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h +++ b/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h @@ -23,9 +23,7 @@ namespace clang::CIRGen { class CIRGenFunction; /// A type-only list of OpenMP clause AST node types. -/// Note: The clause AST classes do not have a default constructor, so a -/// std::tuple is not practical. -template <typename... Clauses> struct OpenMPClauseList {}; +template <typename... Clauses> struct OpenMPNYIClauseList {}; /// Emits OpenMP clauses for a directive, writing results into the /// auto-generated ClauseOps from the OMP dialect. @@ -50,10 +48,11 @@ class OpenMPClauseEmitter { llvm::SmallVectorImpl<const VarDecl *> *mapSyms = nullptr) const; /// Verify the clauses of a directive to make sure all legal cases are either - /// implemented or give a NYI error. If the clause is neither, then - /// an unknown clause error will be emitted. + /// 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. template <typename... SupportedClauses, typename... NYIClauses> - void emitNYI(OpenMPClauseList<NYIClauses...> nyi, + void emitNYI(OpenMPNYIClauseList<NYIClauses...> nyi, llvm::omp::Directive directive) const; private: @@ -63,7 +62,7 @@ class OpenMPClauseEmitter { }; template <typename... SupportedClauses, typename... NYIClauses> -void OpenMPClauseEmitter::emitNYI(OpenMPClauseList<NYIClauses...>, +void OpenMPClauseEmitter::emitNYI(OpenMPNYIClauseList<NYIClauses...>, llvm::omp::Directive directive) const { static_assert( (!isAnyOf<NYIClauses, SupportedClauses...> && ...), @@ -78,7 +77,7 @@ void OpenMPClauseEmitter::emitNYI(OpenMPClauseList<NYIClauses...>, .str(); cgm.errorNYI(c->getBeginLoc(), msg); } else if (!(isa<SupportedClauses>(c) || ...)) { - // Unknown/illegal clause encountered + // Unknown/illegal clause encountered. llvm_unreachable("unexpected OpenMP clause"); } } diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp index 38a55fbdb8c2a..3388f2a9dd3e3 100644 --- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp @@ -40,7 +40,7 @@ CIRGenFunction::emitOMPParallelDirective(const OMPParallelDirective &s) { OpenMPClauseEmitter ce(*this, getCIRGenModule(), builder, begin, s.clauses()); ce.emitProcBind(clauseOps); ce.emitNYI</*supported=*/OMPProcBindClause>( - /*nyi=*/OpenMPClauseList< + /*nyi=*/OpenMPNYIClauseList< OMPAllocateClause, OMPCopyinClause, OMPDefaultClause, OMPFirstprivateClause, OMPIfClause, OMPNumThreadsClause, OMPPrivateClause, OMPReductionClause, OMPSharedClause>{}, @@ -299,7 +299,7 @@ CIRGenFunction::emitOMPTargetDirective(const OMPTargetDirective &s) { OpenMPClauseEmitter ce(*this, getCIRGenModule(), builder, begin, s.clauses()); ce.emitMap(clauseOps, &mapSyms); ce.emitNYI</*supported=*/OMPMapClause>( - /*nyi=*/OpenMPClauseList< + /*nyi=*/OpenMPNYIClauseList< OMPAllocateClause, OMPDefaultClause, OMPDefaultmapClause, OMPDependClause, OMPDeviceClause, OMPFirstprivateClause, OMPHasDeviceAddrClause, OMPIfClause, OMPInReductionClause, >From c6a07032c0f52ec67e1470dc497506f7b690c6a7 Mon Sep 17 00:00:00 2001 From: Jan Leyonberg <[email protected]> Date: Mon, 15 Jun 2026 19:51:20 -0400 Subject: [PATCH 6/6] Fix testx --- clang/test/CIR/CodeGenOpenMP/target-map.c | 20 ++++++++++---------- 1 file changed, 10 insertions(+), 10 deletions(-) diff --git a/clang/test/CIR/CodeGenOpenMP/target-map.c b/clang/test/CIR/CodeGenOpenMP/target-map.c index 22d0ee811b91a..6acbc40f7276a 100644 --- a/clang/test/CIR/CodeGenOpenMP/target-map.c +++ b/clang/test/CIR/CodeGenOpenMP/target-map.c @@ -11,7 +11,7 @@ void use(int); void target_map_to(int x) { // CIR-HOST: cir.func{{.*}}@target_map_to - // CIR-HOST: %[[X_ALLOCA:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["x", init] + // CIR-HOST: %[[X_ALLOCA:.*]] = cir.alloca "x" align(4) init : !cir.ptr<!s32i> // CIR-HOST: %[[MAP:.*]] = omp.map.info var_ptr(%[[X_ALLOCA]] : !cir.ptr<!s32i>, !s32i) map_clauses(to) capture(ByRef) -> !cir.ptr<!s32i> {name = "x"} // CIR-HOST-NEXT: omp.target map_entries(%[[MAP]] -> %[[ARG:.*]] : !cir.ptr<!s32i>) { // CIR-HOST-NEXT: %[[LOAD:.*]] = cir.load align(4) %[[ARG]] @@ -20,7 +20,7 @@ void target_map_to(int x) { // CIR-HOST-NEXT: } // CIR-DEVICE: cir.func{{.*}}@target_map_to - // CIR-DEVICE: %[[X_ALLOCA:.*]] = cir.alloca !s32i, !cir.ptr<!s32i, target_address_space(5)>, ["x", init] + // CIR-DEVICE: %[[X_ALLOCA:.*]] = cir.alloca "x" align(4) init : !cir.ptr<!s32i, target_address_space(5)> // CIR-DEVICE: %[[CAST:.*]] = cir.cast address_space %[[X_ALLOCA]] : !cir.ptr<!s32i, target_address_space(5)> -> !cir.ptr<!s32i> // CIR-DEVICE: %[[MAP:.*]] = omp.map.info var_ptr(%[[CAST]] : !cir.ptr<!s32i>, !s32i) map_clauses(to) capture(ByRef) -> !cir.ptr<!s32i> {name = "x"} // CIR-DEVICE-NEXT: omp.target map_entries(%[[MAP]] -> %[[ARG:.*]] : !cir.ptr<!s32i>) { @@ -34,7 +34,7 @@ void target_map_to(int x) { void target_map_from(int x) { // CIR-HOST: cir.func{{.*}}@target_map_from - // CIR-HOST: %[[X_ALLOCA:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["x", init] + // CIR-HOST: %[[X_ALLOCA:.*]] = cir.alloca "x" align(4) init : !cir.ptr<!s32i> // CIR-HOST: %[[MAP:.*]] = omp.map.info var_ptr(%[[X_ALLOCA]] : !cir.ptr<!s32i>, !s32i) map_clauses(from) capture(ByRef) -> !cir.ptr<!s32i> {name = "x"} // CIR-HOST-NEXT: omp.target map_entries(%[[MAP]] -> %[[ARG:.*]] : !cir.ptr<!s32i>) { // CIR-HOST-NEXT: %[[C42:.*]] = cir.const #cir.int<42> : !s32i @@ -43,7 +43,7 @@ void target_map_from(int x) { // CIR-HOST-NEXT: } // CIR-DEVICE: cir.func{{.*}}@target_map_from - // CIR-DEVICE: %[[X_ALLOCA:.*]] = cir.alloca !s32i, !cir.ptr<!s32i, target_address_space(5)>, ["x", init] + // CIR-DEVICE: %[[X_ALLOCA:.*]] = cir.alloca "x" align(4) init : !cir.ptr<!s32i, target_address_space(5)> // CIR-DEVICE: %[[CAST:.*]] = cir.cast address_space %[[X_ALLOCA]] : !cir.ptr<!s32i, target_address_space(5)> -> !cir.ptr<!s32i> // CIR-DEVICE: %[[MAP:.*]] = omp.map.info var_ptr(%[[CAST]] : !cir.ptr<!s32i>, !s32i) map_clauses(from) capture(ByRef) -> !cir.ptr<!s32i> {name = "x"} // CIR-DEVICE-NEXT: omp.target map_entries(%[[MAP]] -> %[[ARG:.*]] : !cir.ptr<!s32i>) { @@ -57,14 +57,14 @@ void target_map_from(int x) { void target_map_tofrom(int x) { // CIR-HOST: cir.func{{.*}}@target_map_tofrom - // CIR-HOST: %[[X_ALLOCA:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["x", init] + // CIR-HOST: %[[X_ALLOCA:.*]] = cir.alloca "x" align(4) init : !cir.ptr<!s32i> // CIR-HOST: %[[MAP:.*]] = omp.map.info var_ptr(%[[X_ALLOCA]] : !cir.ptr<!s32i>, !s32i) map_clauses(tofrom) capture(ByRef) -> !cir.ptr<!s32i> {name = "x"} // CIR-HOST-NEXT: omp.target map_entries(%[[MAP]] -> %[[ARG:.*]] : !cir.ptr<!s32i>) { // CIR-HOST: omp.terminator // CIR-HOST-NEXT: } // CIR-DEVICE: cir.func{{.*}}@target_map_tofrom - // CIR-DEVICE: %[[X_ALLOCA:.*]] = cir.alloca !s32i, !cir.ptr<!s32i, target_address_space(5)>, ["x", init] + // CIR-DEVICE: %[[X_ALLOCA:.*]] = cir.alloca "x" align(4) init : !cir.ptr<!s32i, target_address_space(5)> // CIR-DEVICE: %[[CAST:.*]] = cir.cast address_space %[[X_ALLOCA]] : !cir.ptr<!s32i, target_address_space(5)> -> !cir.ptr<!s32i> // CIR-DEVICE: %[[MAP:.*]] = omp.map.info var_ptr(%[[CAST]] : !cir.ptr<!s32i>, !s32i) map_clauses(tofrom) capture(ByRef) -> !cir.ptr<!s32i> {name = "x"} // CIR-DEVICE-NEXT: omp.target map_entries(%[[MAP]] -> %[[ARG:.*]] : !cir.ptr<!s32i>) { @@ -78,8 +78,8 @@ void target_map_tofrom(int x) { void target_map_multiple(int a, int b) { // CIR-HOST: cir.func{{.*}}@target_map_multiple - // CIR-HOST-DAG: %[[A_ALLOCA:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["a", init] - // CIR-HOST-DAG: %[[B_ALLOCA:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["b", init] + // CIR-HOST-DAG: %[[A_ALLOCA:.*]] = cir.alloca "a" align(4) init : !cir.ptr<!s32i> + // CIR-HOST-DAG: %[[B_ALLOCA:.*]] = cir.alloca "b" align(4) init : !cir.ptr<!s32i> // CIR-HOST: %[[MAP_A:.*]] = omp.map.info var_ptr(%[[A_ALLOCA]] : !cir.ptr<!s32i>, !s32i) map_clauses(to) capture(ByRef) -> !cir.ptr<!s32i> {name = "a"} // CIR-HOST-NEXT: %[[MAP_B:.*]] = omp.map.info var_ptr(%[[B_ALLOCA]] : !cir.ptr<!s32i>, !s32i) map_clauses(from) capture(ByRef) -> !cir.ptr<!s32i> {name = "b"} // CIR-HOST-NEXT: omp.target map_entries(%[[MAP_A]] -> %[[ARG_A:.*]], %[[MAP_B]] -> %[[ARG_B:.*]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) { @@ -87,8 +87,8 @@ void target_map_multiple(int a, int b) { // CIR-HOST-NEXT: } // CIR-DEVICE: cir.func{{.*}}@target_map_multiple - // CIR-DEVICE-DAG: %[[A_ALLOCA:.*]] = cir.alloca !s32i, !cir.ptr<!s32i, target_address_space(5)>, ["a", init] - // CIR-DEVICE-DAG: %[[B_ALLOCA:.*]] = cir.alloca !s32i, !cir.ptr<!s32i, target_address_space(5)>, ["b", init] + // CIR-DEVICE-DAG: %[[A_ALLOCA:.*]] = cir.alloca "a" align(4) init : !cir.ptr<!s32i, target_address_space(5)> + // CIR-DEVICE-DAG: %[[B_ALLOCA:.*]] = cir.alloca "b" align(4) init : !cir.ptr<!s32i, target_address_space(5)> // CIR-DEVICE: %[[CAST_A:.*]] = cir.cast address_space %[[A_ALLOCA]] : !cir.ptr<!s32i, target_address_space(5)> -> !cir.ptr<!s32i> // CIR-DEVICE: %[[MAP_A:.*]] = omp.map.info var_ptr(%[[CAST_A]] : !cir.ptr<!s32i>, !s32i) map_clauses(to) capture(ByRef) -> !cir.ptr<!s32i> {name = "a"} // CIR-DEVICE: %[[CAST_B:.*]] = cir.cast address_space %[[B_ALLOCA]] : !cir.ptr<!s32i, target_address_space(5)> -> !cir.ptr<!s32i> _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
