Classes | |
| struct | DeviceCodeSet |
| The device code of a module, as gathered by collectDeviceCode. More... | |
Functions | |
| mlir::gpu::GPUModuleOp | getOrCreateGPUModule (mlir::ModuleOp mod, mlir::SymbolTable &symTab) |
Retrieve or create the CUDA Fortran GPU module in the given mod. | |
| bool | isCUDADeviceContext (mlir::Operation *op) |
| bool | isCUDADeviceContext (mlir::Region &, bool isDoConcurrentOffloadEnabled=false) |
| bool | isExecutingOnDevice (mlir::Operation *op) |
| bool | isRegisteredDeviceGlobal (fir::GlobalOp op) |
| bool | isRegisteredDeviceAttr (std::optional< cuf::DataAttribute > attr) |
| bool | isTypeInfoGlobal (fir::GlobalOp op) |
| bool | isTypeDescriptorGlobal (fir::GlobalOp op) |
| bool | isDeviceProcedure (mlir::func::FuncOp funcOp) |
| DeviceCodeSet | collectDeviceCode (mlir::ModuleOp mod, mlir::SymbolTable &symTab, bool rejectDynamicDispatch=false) |
| void | remapProcedureSymbols (mlir::Operation *root, const llvm::DenseMap< mlir::StringAttr, mlir::FlatSymbolRefAttr > &map) |
| void | setDeviceAllocationPolicy (mlir::Operation *func) |
| void | genPointerSync (const mlir::Value box, fir::FirOpBuilder &builder) |
| int | computeElementByteSize (mlir::Location loc, mlir::Type type, fir::KindMapping &kindMap, bool emitErrorOnFailure=true) |
| mlir::Value | computeElementCount (mlir::PatternRewriter &rewriter, mlir::Location loc, mlir::Value shapeOperand, mlir::Type seqType, mlir::Type targetType) |
| void | setDeviceCopyOf (mlir::Operation *copy, llvm::StringRef original) |
Mark copy as the device copy of the procedure named original. | |
| std::optional< llvm::StringRef > | getDeviceCopyOf (mlir::Operation *op) |
The name of the procedure op is the device copy of, if it is one. | |
| cuf::DataAttributeAttr | getDataAttribute (mlir::MLIRContext *mlirContext, std::optional< Fortran::common::CUDADataAttr > cudaAttr) |
| cuf::ProcAttributeAttr | getProcAttribute (mlir::MLIRContext *mlirContext, std::optional< Fortran::common::CUDASubprogramAttrs > cudaAttr) |
| cuf::DataAttributeAttr | getDataAttr (mlir::Operation *op) |
| Returns the data attribute if the operation has one. | |
| bool | hasDataAttr (mlir::Operation *op, cuf::DataAttribute value) |
| Returns true if the operation has a data attribute with the given value. | |
| bool | isDeviceDataAttribute (cuf::DataAttribute attr) |
| bool | hasDeviceDataAttr (mlir::Operation *op) |
| bool | isManagedOrUnifiedDataAttribute (cuf::DataAttribute attr) |
| bool | hasManagedOrUnifiedDataAttr (mlir::Operation *op) |
| void | registerCUFDialectTranslation (mlir::DialectRegistry ®istry) |
| void | populateCUFAllocationConversionPatterns (const fir::LLVMTypeConverter &converter, mlir::DataLayout &dl, const mlir::SymbolTable &symtab, mlir::RewritePatternSet &patterns, llvm::StringRef descriptorAllocFunction={}, llvm::StringRef descriptorFreeFunction={}) |
| void | populateCUFGPUToLLVMConversionPatterns (fir::LLVMTypeConverter &converter, mlir::RewritePatternSet &patterns, mlir::PatternBenefit benefit=1) |
| void | populateCUFToFIRConversionPatterns (const fir::LLVMTypeConverter &converter, mlir::DataLayout &dl, const mlir::SymbolTable &symtab, mlir::RewritePatternSet &patterns) |
| Patterns that convert CUF operations to runtime calls. | |
| void | populateFIRCUFConversionPatterns (const mlir::SymbolTable &symtab, mlir::RewritePatternSet &patterns) |
| Patterns that updates fir operations in presence of CUF. | |
Detailed Description
Instantiation of pft::Variable in FIR/MLIR.
Function Documentation
◆ collectDeviceCode()
| cuf::DeviceCodeSet cuf::collectDeviceCode | ( | mlir::ModuleOp | mod, |
| mlir::SymbolTable & | symTab, | ||
| bool | rejectDynamicDispatch = false ) |
Gather the device code of mod. A host_device procedure that already has a device copy is host code and is not walked for callees. With rejectDynamicDispatch, a type-bound call with dynamic dispatch in device code is reported as not yet implemented.
◆ getOrCreateGPUModule()
| mlir::gpu::GPUModuleOp cuf::getOrCreateGPUModule | ( | mlir::ModuleOp | mod, |
| mlir::SymbolTable & | symTab ) |
◆ hasDeviceDataAttr()
| bool cuf::hasDeviceDataAttr | ( | mlir::Operation * | op | ) |
Returns true if the operation has a cuf::DataAttributeAttr whose value returns true for isDeviceDataAttribute.
◆ hasManagedOrUnifiedDataAttr()
| bool cuf::hasManagedOrUnifiedDataAttr | ( | mlir::Operation * | op | ) |
Returns true if the operation has a cuf::DataAttributeAttr whose value returns true for isManagedOrUnifiedDataAttribute.
◆ isDeviceDataAttribute()
| bool cuf::isDeviceDataAttribute | ( | cuf::DataAttribute | attr | ) |
Check if a CUDA data attribute represents device data. Returns true for Device, Managed, Constant, Shared, and Unified attributes.
◆ isDeviceProcedure()
| bool cuf::isDeviceProcedure | ( | mlir::func::FuncOp | funcOp | ) |
True for procedures that have a device side: attributes(device), (global), (grid_global) and (host,device). Unlike isCUDADeviceContext, host_device counts, since its body is compiled for the device as well.
◆ isExecutingOnDevice()
| bool cuf::isExecutingOnDevice | ( | mlir::Operation * | op | ) |
True when op is compiled for device execution: a CUDA device procedure, a GPU module or launch, or an OpenACC offload region. Host and host_device procedures return false. Unlike isCUDADeviceContext, gpu.module and gpu.launch count even without a CUDA procedure attribute, and host_data does not.
◆ isManagedOrUnifiedDataAttribute()
| bool cuf::isManagedOrUnifiedDataAttribute | ( | cuf::DataAttribute | attr | ) |
Check if a CUDA data attribute represents managed or unified storage. Returns true for Managed and Unified attributes. Such storage is device-accessible but may migrate rather than being statically guaranteed to reside on the device.
◆ isTypeDescriptorGlobal()
| bool cuf::isTypeDescriptorGlobal | ( | fir::GlobalOp | op | ) |
True when op is the type descriptor of a derived type, the type-info global whose address is stored in descriptors.
◆ isTypeInfoGlobal()
| bool cuf::isTypeInfoGlobal | ( | fir::GlobalOp | op | ) |
True when op holds runtime type information generated for a derived type (type descriptor, binding table, component table, names, ...).
◆ populateCUFAllocationConversionPatterns()
| void cuf::populateCUFAllocationConversionPatterns | ( | const fir::LLVMTypeConverter & | converter, |
| mlir::DataLayout & | dl, | ||
| const mlir::SymbolTable & | symtab, | ||
| mlir::RewritePatternSet & | patterns, | ||
| llvm::StringRef | descriptorAllocFunction = {}, | ||
| llvm::StringRef | descriptorFreeFunction = {} ) |
Patterns that convert CUF operations to runtime calls. descriptorAllocFunction / descriptorFreeFunction, when non-empty, override the runtime functions used for descriptor allocations / frees (same signatures as CUFAllocDescriptor / CUFFreeDescriptor).
◆ registerCUFDialectTranslation()
| void cuf::registerCUFDialectTranslation | ( | mlir::DialectRegistry & | registry | ) |
Register the CUF dialect and the translation from it to the LLVM IR in the given registry.
◆ remapProcedureSymbols()
| void cuf::remapProcedureSymbols | ( | mlir::Operation * | root, |
| const llvm::DenseMap< mlir::StringAttr, mlir::FlatSymbolRefAttr > & | map ) |
Point every fir.call and fir.address_of under root whose symbol is a key of map at the mapped symbol instead.
◆ setDeviceAllocationPolicy()
| void cuf::setDeviceAllocationPolicy | ( | mlir::Operation * | func | ) |
Record on func the allocation policy of device code: the policy in effect for it with stack arrays disabled, since the device stack is far smaller than the host one.
Generated on for FLANG by
1.14.0