16#include "mlir/Dialect/OpenMP/OpenMPDialect.h"
17#include "mlir/Dialect/OpenMP/OpenMPInterfaces.h"
23static mlir::omp::DeclareTargetDeviceType
26 case OMPDeclareTargetDeclAttr::DT_Host:
27 return mlir::omp::DeclareTargetDeviceType::host;
28 case OMPDeclareTargetDeclAttr::DT_NoHost:
29 return mlir::omp::DeclareTargetDeviceType::nohost;
30 case OMPDeclareTargetDeclAttr::DT_Any:
31 return mlir::omp::DeclareTargetDeviceType::any;
33 llvm_unreachable(
"unexpected device type");
36static mlir::omp::DeclareTargetCaptureClause
39 case OMPDeclareTargetDeclAttr::MT_To:
40 return mlir::omp::DeclareTargetCaptureClause::to;
41 case OMPDeclareTargetDeclAttr::MT_Enter:
42 return mlir::omp::DeclareTargetCaptureClause::enter;
43 case OMPDeclareTargetDeclAttr::MT_Link:
44 return mlir::omp::DeclareTargetCaptureClause::link;
45 case OMPDeclareTargetDeclAttr::MT_Local:
46 return mlir::omp::DeclareTargetCaptureClause::none;
48 llvm_unreachable(
"unexpected map type");
54 std::optional<OMPDeclareTargetDeclAttr::DevTypeTy> devTy =
55 OMPDeclareTargetDeclAttr::getDeviceType(vd);
59 if (!isDevice && *devTy == OMPDeclareTargetDeclAttr::DT_NoHost)
62 if (isDevice && *devTy == OMPDeclareTargetDeclAttr::DT_Host)
74 if (
const auto *e = dyn_cast<OMPExecutableDirective>(
s))
77 for (
const Stmt *child :
s->children())
83bool CIRGenOpenMPRuntime::emitTargetFunctions(
GlobalDecl gd) {
84 bool isDevice = cgm.getLangOpts().OpenMPIsTargetDevice;
87 if (
const auto *fd = dyn_cast<FunctionDecl>(gd.
getDecl()))
95 if (
const auto *fd = dyn_cast<FunctionDecl>(vd))
100 if (OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(vd) ||
101 alreadyEmittedTargetDecls.count(vd) != 0)
107 if (
const auto *fd = dyn_cast<FunctionDecl>(vd))
108 if (fd->doesThisDeclarationHaveABody() &&
115bool CIRGenOpenMPRuntime::emitTargetGlobalVariable(GlobalDecl gd) {
117 cgm.getLangOpts().OpenMPIsTargetDevice))
120 if (!cgm.getLangOpts().OpenMPIsTargetDevice)
131 std::optional<OMPDeclareTargetDeclAttr::MapTypeTy> res =
132 OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(vd);
133 if (!res || *res == OMPDeclareTargetDeclAttr::MT_Link ||
134 ((*res == OMPDeclareTargetDeclAttr::MT_To ||
135 *res == OMPDeclareTargetDeclAttr::MT_Enter) &&
137 if (res && *res == OMPDeclareTargetDeclAttr::MT_Link)
138 cgm.errorNYI(vd->getSourceRange(),
139 "declare target global variable with link clause");
150 return emitTargetFunctions(gd);
152 return emitTargetGlobalVariable(gd);
156 if (!cgm.getLangOpts().OpenMPIsTargetDevice)
161 if (OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(d)) {
162 if (d->hasBody() && alreadyEmittedTargetDecls.count(d) == 0) {
163 auto f = dyn_cast_if_present<cir::FuncOp>(
164 cgm.getGlobalValue(cgm.getMangledName(gd)));
166 return !f.isDeclaration();
172 return !alreadyEmittedTargetDecls.insert(d).second;
176 cir::FuncOp funcOp) {
177 const auto *
attr = fd->
getAttr<OMPDeclareTargetDeclAttr>();
178 assert(
attr &&
"expected OMPDeclareTargetDeclAttr");
183 if (std::optional<OMPDeclareTargetDeclAttr *> activeAttr =
184 OMPDeclareTargetDeclAttr::getActiveAttr(fd))
185 if ((*activeAttr)->getIndirect())
187 "declare target function with indirect clause");
189 auto declTargetIface =
190 llvm::cast<mlir::omp::DeclareTargetInterface>(funcOp.getOperation());
static mlir::omp::DeclareTargetCaptureClause convertCaptureClause(OMPDeclareTargetDeclAttr::MapTypeTy mapTy)
static bool containsTargetRegion(const Stmt *s)
Recursively check whether the statement tree contains any OpenMP target execution directive (e....
static bool isAssumedToBeNotEmitted(const ValueDecl *vd, bool isDevice)
Returns true if the declaration should be skipped based on its device_type attribute and the current ...
static mlir::omp::DeclareTargetDeviceType convertDeviceType(OMPDeclareTargetDeclAttr::DevTypeTy devTy)
This file defines OpenMP AST classes for clauses.
__device__ __2f16 float __ockl_bool s
void emitDeclareTargetFunction(const FunctionDecl *fd, cir::FuncOp funcOp)
If the function has an OMPDeclareTargetDeclAttr, set the corresponding omp.declare_target attribute o...
bool emitTargetGlobal(GlobalDecl gd)
Check whether the given GlobalDecl needs special handling for device compilation.
bool markAsGlobalTarget(GlobalDecl gd)
Mark a function reference as one that should be emitted on the device.
Represents a function declaration or definition.
SourceRange getSourceRange() const override LLVM_READONLY
Source range that this declaration covers.
GlobalDecl - represents a global declaration.
const Decl * getDecl() const
Stmt - This represents one statement.
Represent the declaration of a variable (in which case it is an lvalue) a function (in which case it ...
The JSON file list parser is used to communicate input to InstallAPI.
bool isa(CodeGen::Address addr)
bool isOpenMPTargetExecutionDirective(OpenMPDirectiveKind DKind)
Checks if the specified directive is a target code offload directive.
U cast(CodeGen::Address addr)