158 lines · cpp
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-exception6//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/Dialect/Support/KindMapping.h"13#include "flang/Optimizer/HLFIR/HLFIROps.h"14#include "mlir/Dialect/Func/IR/FuncOps.h"15#include "mlir/Dialect/LLVMIR/NVVMDialect.h"16#include "mlir/Dialect/OpenACC/OpenACC.h"17 18/// Retrieve or create the CUDA Fortran GPU module in the give in \p mod.19mlir::gpu::GPUModuleOp cuf::getOrCreateGPUModule(mlir::ModuleOp mod,20 mlir::SymbolTable &symTab) {21 if (auto gpuMod = symTab.lookup<mlir::gpu::GPUModuleOp>(cudaDeviceModuleName))22 return gpuMod;23 24 auto *ctx = mod.getContext();25 mod->setAttr(mlir::gpu::GPUDialect::getContainerModuleAttrName(),26 mlir::UnitAttr::get(ctx));27 28 mlir::OpBuilder builder(ctx);29 auto gpuMod = mlir::gpu::GPUModuleOp::create(builder, mod.getLoc(),30 cudaDeviceModuleName);31 mlir::Block::iterator insertPt(mod.getBodyRegion().front().end());32 symTab.insert(gpuMod, insertPt);33 return gpuMod;34}35 36bool cuf::isCUDADeviceContext(mlir::Operation *op) {37 if (!op || !op->getParentRegion())38 return false;39 return isCUDADeviceContext(*op->getParentRegion());40}41 42// Check if the insertion point is currently in a device context. HostDevice43// subprogram are not considered fully device context so it will return false44// for it.45// If the insertion point is inside an OpenACC region op, it is considered46// device context.47bool cuf::isCUDADeviceContext(mlir::Region ®ion,48 bool isDoConcurrentOffloadEnabled) {49 if (region.getParentOfType<cuf::KernelOp>())50 return true;51 if (region.getParentOfType<mlir::acc::ComputeRegionOpInterface>())52 return true;53 if (auto funcOp = region.getParentOfType<mlir::func::FuncOp>()) {54 if (auto cudaProcAttr =55 funcOp.getOperation()->getAttrOfType<cuf::ProcAttributeAttr>(56 cuf::getProcAttrName())) {57 return cudaProcAttr.getValue() != cuf::ProcAttribute::Host &&58 cudaProcAttr.getValue() != cuf::ProcAttribute::HostDevice;59 }60 }61 if (isDoConcurrentOffloadEnabled &&62 region.getParentOfType<fir::DoConcurrentLoopOp>())63 return true;64 return false;65}66 67bool cuf::isRegisteredDeviceAttr(std::optional<cuf::DataAttribute> attr) {68 if (attr && (*attr == cuf::DataAttribute::Device ||69 *attr == cuf::DataAttribute::Managed ||70 *attr == cuf::DataAttribute::Constant))71 return true;72 return false;73}74 75bool cuf::isRegisteredDeviceGlobal(fir::GlobalOp op) {76 if (op.getConstant())77 return false;78 return isRegisteredDeviceAttr(op.getDataAttr());79}80 81void cuf::genPointerSync(const mlir::Value box, fir::FirOpBuilder &builder) {82 if (auto declareOp = box.getDefiningOp<hlfir::DeclareOp>()) {83 if (auto addrOfOp = declareOp.getMemref().getDefiningOp<fir::AddrOfOp>()) {84 auto mod = addrOfOp->getParentOfType<mlir::ModuleOp>();85 if (auto globalOp =86 mod.lookupSymbol<fir::GlobalOp>(addrOfOp.getSymbol())) {87 if (cuf::isRegisteredDeviceGlobal(globalOp)) {88 cuf::SyncDescriptorOp::create(builder, box.getLoc(),89 addrOfOp.getSymbol());90 }91 }92 }93 }94}95 96int cuf::computeElementByteSize(mlir::Location loc, mlir::Type type,97 fir::KindMapping &kindMap,98 bool emitErrorOnFailure) {99 auto eleTy = fir::unwrapSequenceType(type);100 if (auto t{mlir::dyn_cast<mlir::IntegerType>(eleTy)})101 return t.getWidth() / 8;102 if (auto t{mlir::dyn_cast<mlir::FloatType>(eleTy)})103 return t.getWidth() / 8;104 if (auto t{mlir::dyn_cast<fir::LogicalType>(eleTy)})105 return kindMap.getLogicalBitsize(t.getFKind()) / 8;106 if (auto t{mlir::dyn_cast<mlir::ComplexType>(eleTy)}) {107 int elemSize =108 mlir::cast<mlir::FloatType>(t.getElementType()).getWidth() / 8;109 return 2 * elemSize;110 }111 if (auto t{mlir::dyn_cast<fir::CharacterType>(eleTy)})112 return kindMap.getCharacterBitsize(t.getFKind()) / 8;113 if (emitErrorOnFailure)114 mlir::emitError(loc, "unsupported type");115 return 0;116}117 118mlir::Value cuf::computeElementCount(mlir::PatternRewriter &rewriter,119 mlir::Location loc,120 mlir::Value shapeOperand,121 mlir::Type seqType,122 mlir::Type targetType) {123 if (shapeOperand) {124 // Dynamic extent - extract from shape operand125 llvm::SmallVector<mlir::Value> extents;126 if (auto shapeOp =127 mlir::dyn_cast<fir::ShapeOp>(shapeOperand.getDefiningOp())) {128 extents = shapeOp.getExtents();129 } else if (auto shapeShiftOp = mlir::dyn_cast<fir::ShapeShiftOp>(130 shapeOperand.getDefiningOp())) {131 for (auto i : llvm::enumerate(shapeShiftOp.getPairs()))132 if (i.index() & 1)133 extents.push_back(i.value());134 }135 136 if (extents.empty())137 return mlir::Value();138 139 // Compute total element count by multiplying all dimensions140 mlir::Value count =141 fir::ConvertOp::create(rewriter, loc, targetType, extents[0]);142 for (unsigned i = 1; i < extents.size(); ++i) {143 auto operand =144 fir::ConvertOp::create(rewriter, loc, targetType, extents[i]);145 count = mlir::arith::MulIOp::create(rewriter, loc, count, operand);146 }147 return count;148 } else {149 // Static extent - use constant array size150 if (auto seqTy = mlir::dyn_cast_or_null<fir::SequenceType>(seqType)) {151 mlir::IntegerAttr attr =152 rewriter.getIntegerAttr(targetType, seqTy.getConstantArraySize());153 return mlir::arith::ConstantOp::create(rewriter, loc, targetType, attr);154 }155 }156 return mlir::Value();157}158