1//===-- CUFCommon.cpp - Shared functions between passes ---------*- C++ -*-===//
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#include "flang/Optimizer/Builder/CUFCommon.h"
10#include "flang/Optimizer/Builder/FIRBuilder.h"
11#include "flang/Optimizer/Dialect/CUF/CUFOps.h"
12#include "flang/Optimizer/HLFIR/HLFIROps.h"
13#include "mlir/Dialect/Func/IR/FuncOps.h"
14#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
15#include "mlir/Dialect/OpenACC/OpenACC.h"
16
17/// Retrieve or create the CUDA Fortran GPU module in the give in \p mod.
18mlir::gpu::GPUModuleOp cuf::getOrCreateGPUModule(mlir::ModuleOp mod,
19 mlir::SymbolTable &symTab) {
20 if (auto gpuMod = symTab.lookup<mlir::gpu::GPUModuleOp>(cudaDeviceModuleName))
21 return gpuMod;
22
23 auto *ctx = mod.getContext();
24 mod->setAttr(mlir::gpu::GPUDialect::getContainerModuleAttrName(),
25 mlir::UnitAttr::get(ctx));
26
27 mlir::OpBuilder builder(ctx);
28 auto gpuMod = builder.create<mlir::gpu::GPUModuleOp>(mod.getLoc(),
29 cudaDeviceModuleName);
30 mlir::Block::iterator insertPt(mod.getBodyRegion().front().end());
31 symTab.insert(gpuMod, insertPt);
32 return gpuMod;
33}
34
35bool cuf::isCUDADeviceContext(mlir::Operation *op) {
36 if (!op || !op->getParentRegion())
37 return false;
38 return isCUDADeviceContext(*op->getParentRegion());
39}
40
41// Check if the insertion point is currently in a device context. HostDevice
42// subprogram are not considered fully device context so it will return false
43// for it.
44// If the insertion point is inside an OpenACC region op, it is considered
45// device context.
46bool cuf::isCUDADeviceContext(mlir::Region &region,
47 bool isDoConcurrentOffloadEnabled) {
48 if (region.getParentOfType<cuf::KernelOp>())
49 return true;
50 if (region.getParentOfType<mlir::acc::ComputeRegionOpInterface>())
51 return true;
52 if (auto funcOp = region.getParentOfType<mlir::func::FuncOp>()) {
53 if (auto cudaProcAttr =
54 funcOp.getOperation()->getAttrOfType<cuf::ProcAttributeAttr>(
55 cuf::getProcAttrName())) {
56 return cudaProcAttr.getValue() != cuf::ProcAttribute::Host &&
57 cudaProcAttr.getValue() != cuf::ProcAttribute::HostDevice;
58 }
59 }
60 if (isDoConcurrentOffloadEnabled &&
61 region.getParentOfType<fir::DoConcurrentLoopOp>())
62 return true;
63 return false;
64}
65
66bool cuf::isRegisteredDeviceAttr(std::optional<cuf::DataAttribute> attr) {
67 if (attr && (*attr == cuf::DataAttribute::Device ||
68 *attr == cuf::DataAttribute::Managed ||
69 *attr == cuf::DataAttribute::Constant))
70 return true;
71 return false;
72}
73
74bool cuf::isRegisteredDeviceGlobal(fir::GlobalOp op) {
75 if (op.getConstant())
76 return false;
77 return isRegisteredDeviceAttr(op.getDataAttr());
78}
79
80void cuf::genPointerSync(const mlir::Value box, fir::FirOpBuilder &builder) {
81 if (auto declareOp = box.getDefiningOp<hlfir::DeclareOp>()) {
82 if (auto addrOfOp = declareOp.getMemref().getDefiningOp<fir::AddrOfOp>()) {
83 auto mod = addrOfOp->getParentOfType<mlir::ModuleOp>();
84 if (auto globalOp =
85 mod.lookupSymbol<fir::GlobalOp>(addrOfOp.getSymbol())) {
86 if (cuf::isRegisteredDeviceGlobal(globalOp)) {
87 builder.create<cuf::SyncDescriptorOp>(box.getLoc(),
88 addrOfOp.getSymbol());
89 }
90 }
91 }
92 }
93}
94

source code of flang/lib/Optimizer/Builder/CUFCommon.cpp