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 if (region.getParentOfType<cuf::KernelOp>())
48 return true;
49 if (region.getParentOfType<mlir::acc::ComputeRegionOpInterface>())
50 return true;
51 if (auto funcOp = region.getParentOfType<mlir::func::FuncOp>()) {
52 if (auto cudaProcAttr =
53 funcOp.getOperation()->getAttrOfType<cuf::ProcAttributeAttr>(
54 cuf::getProcAttrName())) {
55 return cudaProcAttr.getValue() != cuf::ProcAttribute::Host &&
56 cudaProcAttr.getValue() != cuf::ProcAttribute::HostDevice;
57 }
58 }
59 return false;
60}
61
62bool cuf::isRegisteredDeviceAttr(std::optional<cuf::DataAttribute> attr) {
63 if (attr && (*attr == cuf::DataAttribute::Device ||
64 *attr == cuf::DataAttribute::Managed ||
65 *attr == cuf::DataAttribute::Constant))
66 return true;
67 return false;
68}
69
70bool cuf::isRegisteredDeviceGlobal(fir::GlobalOp op) {
71 if (op.getConstant())
72 return false;
73 return isRegisteredDeviceAttr(op.getDataAttr());
74}
75
76void cuf::genPointerSync(const mlir::Value box, fir::FirOpBuilder &builder) {
77 if (auto declareOp = box.getDefiningOp<hlfir::DeclareOp>()) {
78 if (auto addrOfOp = declareOp.getMemref().getDefiningOp<fir::AddrOfOp>()) {
79 auto mod = addrOfOp->getParentOfType<mlir::ModuleOp>();
80 if (auto globalOp =
81 mod.lookupSymbol<fir::GlobalOp>(addrOfOp.getSymbol())) {
82 if (cuf::isRegisteredDeviceGlobal(globalOp)) {
83 builder.create<cuf::SyncDescriptorOp>(box.getLoc(),
84 addrOfOp.getSymbol());
85 }
86 }
87 }
88 }
89}
90

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