clang 23.0.0git
CIRGenOpenMPRuntime.cpp
Go to the documentation of this file.
1//===--- CIRGenOpenMPRuntime.cpp - OpenMP code generation helpers ------=--===//
2//
3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4// See https://llvm.org/LICENSE.txt for license information.
5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6//
7//===----------------------------------------------------------------------===//
8//
9// Helpers for OpenMP-specific CIR code generation.
10//
11////===--------------------------------------------------------------------===//
12
13#include "CIRGenOpenMPRuntime.h"
14#include "CIRGenModule.h"
15
16#include "mlir/Dialect/OpenMP/OpenMPDialect.h"
17#include "mlir/Dialect/OpenMP/OpenMPInterfaces.h"
19
20using namespace clang;
21using namespace clang::CIRGen;
22
23static mlir::omp::DeclareTargetDeviceType
24convertDeviceType(OMPDeclareTargetDeclAttr::DevTypeTy devTy) {
25 switch (devTy) {
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;
32 }
33 llvm_unreachable("unexpected device type");
34}
35
36static mlir::omp::DeclareTargetCaptureClause
37convertCaptureClause(OMPDeclareTargetDeclAttr::MapTypeTy mapTy) {
38 switch (mapTy) {
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;
47 }
48 llvm_unreachable("unexpected map type");
49}
50
51/// Returns true if the declaration should be skipped based on its
52/// device_type attribute and the current compilation mode.
53static bool isAssumedToBeNotEmitted(const ValueDecl *vd, bool isDevice) {
54 std::optional<OMPDeclareTargetDeclAttr::DevTypeTy> devTy =
55 OMPDeclareTargetDeclAttr::getDeviceType(vd);
56 if (!devTy)
57 return false;
58 // Do not emit device_type(nohost) functions for the host.
59 if (!isDevice && *devTy == OMPDeclareTargetDeclAttr::DT_NoHost)
60 return true;
61 // Do not emit device_type(host) functions for the device.
62 if (isDevice && *devTy == OMPDeclareTargetDeclAttr::DT_Host)
63 return true;
64 return false;
65}
66
67/// Recursively check whether the statement tree contains any OpenMP target
68/// execution directive (e.g. 'omp target', 'omp target parallel', etc.).
69/// Used to identify host functions that must be emitted on the device because
70/// they contain target regions that will be outlined during MLIR lowering.
71static bool containsTargetRegion(const Stmt *s) {
72 if (!s)
73 return false;
74 if (const auto *e = dyn_cast<OMPExecutableDirective>(s))
75 if (isOpenMPTargetExecutionDirective(e->getDirectiveKind()))
76 return true;
77 for (const Stmt *child : s->children())
78 if (containsTargetRegion(child))
79 return true;
80 return false;
81}
82
83bool CIRGenOpenMPRuntime::emitTargetFunctions(GlobalDecl gd) {
84 bool isDevice = cgm.getLangOpts().OpenMPIsTargetDevice;
85
86 if (!isDevice) {
87 if (const auto *fd = dyn_cast<FunctionDecl>(gd.getDecl()))
89 return true;
90 return false;
91 }
92
93 const auto *vd = cast<ValueDecl>(gd.getDecl());
94
95 if (const auto *fd = dyn_cast<FunctionDecl>(vd))
97 return true;
98
99 // Do not emit function if it is not marked as declare target.
100 if (OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(vd) ||
101 alreadyEmittedTargetDecls.count(vd) != 0)
102 return false;
103
104 // We must also host functions that contain target regions,
105 // because the omp.target ops are nested inside the host function rather than
106 // being outlined early. The containsTargetRegion check handles this.
107 if (const auto *fd = dyn_cast<FunctionDecl>(vd))
108 if (fd->doesThisDeclarationHaveABody() &&
109 containsTargetRegion(fd->getBody()))
110 return false;
111
112 return true;
113}
114
115bool CIRGenOpenMPRuntime::emitTargetGlobalVariable(GlobalDecl gd) {
117 cgm.getLangOpts().OpenMPIsTargetDevice))
118 return true;
119
120 if (!cgm.getLangOpts().OpenMPIsTargetDevice)
121 return false;
122
123 // We do not need to scan for target regions since there's not early
124 // outlining like in OGCG, they will be emitted as omp.target ops instead.
125
126 const auto *vd = cast<VarDecl>(gd.getDecl());
127
128 // Do not emit variable if it is not marked as declare target.
129 // OGCG also defers link-clause and USM variables here; we emit errorNYI
130 // for those since they are not yet supported.
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) &&
136 false /* NYI: HasRequiresUnifiedSharedMemory */)) {
137 if (res && *res == OMPDeclareTargetDeclAttr::MT_Link)
138 cgm.errorNYI(vd->getSourceRange(),
139 "declare target global variable with link clause");
140 // OGCG defers these variables for later emission. We skip them for now.
141 return true;
142 }
143 return false;
144}
145
146// Mirrors CGOpenMPRuntime::emitTargetGlobal.
148 if (isa<FunctionDecl>(gd.getDecl()) ||
150 return emitTargetFunctions(gd);
151
152 return emitTargetGlobalVariable(gd);
153}
154
156 if (!cgm.getLangOpts().OpenMPIsTargetDevice)
157 return true;
158
159 const auto *d = cast<FunctionDecl>(gd.getDecl());
160
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)));
165 if (f)
166 return !f.isDeclaration();
167 return false;
168 }
169 return true;
170 }
171
172 return !alreadyEmittedTargetDecls.insert(d).second;
173}
174
176 cir::FuncOp funcOp) {
177 const auto *attr = fd->getAttr<OMPDeclareTargetDeclAttr>();
178 assert(attr && "expected OMPDeclareTargetDeclAttr");
179
180 // Handles the 'indirect' clause here by creating a global variable to hold
181 // the device function address for runtime resolution of indirect calls on
182 // the device.
183 if (std::optional<OMPDeclareTargetDeclAttr *> activeAttr =
184 OMPDeclareTargetDeclAttr::getActiveAttr(fd))
185 if ((*activeAttr)->getIndirect())
186 cgm.errorNYI(fd->getSourceRange(),
187 "declare target function with indirect clause");
188
189 auto declTargetIface =
190 llvm::cast<mlir::omp::DeclareTargetInterface>(funcOp.getOperation());
191 declTargetIface.setDeclareTarget(convertDeviceType(attr->getDevType()),
192 convertCaptureClause(attr->getMapType()),
193 /*automap=*/false);
194}
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.
T * getAttr() const
Definition DeclBase.h:581
Represents a function declaration or definition.
Definition Decl.h:2015
SourceRange getSourceRange() const override LLVM_READONLY
Source range that this declaration covers.
Definition Decl.cpp:4515
GlobalDecl - represents a global declaration.
Definition GlobalDecl.h:57
const Decl * getDecl() const
Definition GlobalDecl.h:106
Stmt - This represents one statement.
Definition Stmt.h:86
Represent the declaration of a variable (in which case it is an lvalue) a function (in which case it ...
Definition Decl.h:712
The JSON file list parser is used to communicate input to InstallAPI.
bool isa(CodeGen::Address addr)
Definition Address.h:330
bool isOpenMPTargetExecutionDirective(OpenMPDirectiveKind DKind)
Checks if the specified directive is a target code offload directive.
U cast(CodeGen::Address addr)
Definition Address.h:327