Author: Jan Leyonberg Date: 2026-06-25T10:49:08-04:00 New Revision: 213c7b7634d8cc585368e5acc0e26002b3715495
URL: https://github.com/llvm/llvm-project/commit/213c7b7634d8cc585368e5acc0e26002b3715495 DIFF: https://github.com/llvm/llvm-project/commit/213c7b7634d8cc585368e5acc0e26002b3715495.diff LOG: [CIR][OpenMP] Initial implementation of target region support (#195452) 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 Added: clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h clang/test/CIR/CodeGenOpenMP/target-map-llvm-device.c clang/test/CIR/CodeGenOpenMP/target-map-llvm-host.c clang/test/CIR/CodeGenOpenMP/target-map.c Modified: clang/lib/CIR/CodeGen/CIRGenFunction.h clang/lib/CIR/CodeGen/CIRGenOpenMPClause.cpp clang/lib/CIR/CodeGen/CIRGenStmt.cpp clang/lib/CIR/CodeGen/CIRGenStmtOpenMP.cpp clang/test/CIR/CodeGenOpenMP/not-yet-implemented.c mlir/include/mlir/Dialect/OpenMP/OpenMPDialect.h mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp Removed: ################################################################################ diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h index b6a4a277fab92..0ce6005ccac88 100644 --- a/clang/lib/CIR/CodeGen/CIRGenFunction.h +++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h @@ -2566,10 +2566,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..16ac4440660b5 100644 --- a/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.cpp @@ -6,90 +6,140 @@ // //===----------------------------------------------------------------------===// // -// Emit OpenMP clause nodes as CIR code. +// OpenMP clause emitter implementation. // //===----------------------------------------------------------------------===// +#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 OpenMPClauseEmitter::emitProcBind( + mlir::omp::ProcBindClauseOps &result) const { + for (const OMPClause *clause : clauses) { + const auto *pbc = dyn_cast<OMPProcBindClause>(clause); + if (!pbc) + continue; + + 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::convertProcBindKind(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 OpenMPClauseEmitter::emitMap( + 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..54c7366b1d769 --- /dev/null +++ b/clang/lib/CIR/CodeGen/CIRGenOpenMPClause.h @@ -0,0 +1,88 @@ +//===--- 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. +// 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" + +#include <type_traits> + +namespace clang::CIRGen { + +class CIRGenFunction; + +/// A type-only list of OpenMP clause AST node types. +template <typename... Clauses> struct OpenMPNYIClauseList {}; + +/// Emits OpenMP clauses for a directive, writing results into the +/// auto-generated ClauseOps from the OMP dialect. +class OpenMPClauseEmitter { + CIRGenFunction &cgf; + CIRGenModule &cgm; + CIRGenBuilderTy &builder; + mlir::Location loc; + llvm::ArrayRef<const OMPClause *> clauses; + +public: + 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 emitProcBind(mlir::omp::ProcBindClauseOps &result) const; + + /// Emit map clauses. The optional \p mapSyms parameter collects the + /// VarDecls corresponding to each map operand. + bool emitMap(mlir::omp::MapClauseOps &result, + 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. 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(OpenMPNYIClauseList<NYIClauses...> nyi, + llvm::omp::Directive directive) const; + +private: + /// 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... SupportedClauses, typename... NYIClauses> +void 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"); + + for (const OMPClause *c : clauses) { + if (isa<NYIClauses...>(c)) { + std::string msg = + (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)) { + // Unknown/illegal clause encountered. + llvm_unreachable("unexpected OpenMP clause"); + } + } +} + +} // 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 47c94cb4ec535..a2c999f584399 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 must 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..17a1fb8090f5c 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,20 @@ 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; + OpenMPClauseEmitter ce(*this, getCIRGenModule(), builder, begin, s.clauses()); + ce.emitProcBind(clauseOps); + ce.emitNYI</*supported=*/OMPProcBindClause>( + /*nyi=*/OpenMPNYIClauseList< + 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 +215,108 @@ 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::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); } 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..29e280bf262d0 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..b6fa078d143c5 --- /dev/null +++ b/clang/test/CIR/CodeGenOpenMP/target-map-llvm-device.c @@ -0,0 +1,130 @@ +// 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-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: %[[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: %[[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: %[[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 %[[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: %[[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 + +// 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..a42e955396980 --- /dev/null +++ b/clang/test/CIR/CodeGenOpenMP/target-map-llvm-host.c @@ -0,0 +1,163 @@ +// 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( +// 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{{.*}}(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{{.*}}(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{{.*}}(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{{.*}}(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-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-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-SAME: (ptr %[[ARG:[^,]+]], ptr +// LLVM: %[[LD:.*]] = load i32, ptr %[[ARG]], align 4 +// LLVM: %[[ADD:.*]] = add nsw i32 %[[LD]], 1 +// LLVM: store i32 %[[ADD]], ptr %[[ARG]], align 4 +// LLVM: ret void + +// LLVM-LABEL: define internal void @__omp_offloading_{{.*}}_target_map_multiple_l +// 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. + +// 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..394b7abb9e0b2 --- /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 "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 kernel_type(generic) 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 "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 kernel_type(generic) 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 "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 kernel_type(generic) 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 "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 kernel_type(generic) 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 "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 kernel_type(generic) 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 "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 kernel_type(generic) 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 "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 kernel_type(generic) 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 "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> + // 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 kernel_type(generic) 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 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 b0d67e73aab30..d29d221dfed14 100644 --- a/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp +++ b/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp @@ -3921,6 +3921,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())); _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
