diff --git a/flang/include/flang/Optimizer/Support/DataLayout.h b/flang/include/flang/Optimizer/Support/DataLayout.h index 6072425b7d63..957ea99162c5 100644 --- a/flang/include/flang/Optimizer/Support/DataLayout.h +++ b/flang/include/flang/Optimizer/Support/DataLayout.h @@ -18,10 +18,14 @@ namespace mlir { class ModuleOp; -} +namespace gpu { +class GPUModuleOp; +} // namespace gpu +} // namespace mlir + namespace llvm { class DataLayout; -} +} // namespace llvm namespace fir::support { /// Create an mlir::DataLayoutSpecInterface attribute from an llvm::DataLayout @@ -30,6 +34,8 @@ namespace fir::support { /// the llvm::DataLayout on the module. /// These attributes are replaced if they were already set. void setMLIRDataLayout(mlir::ModuleOp mlirModule, const llvm::DataLayout &dl); +void setMLIRDataLayout(mlir::gpu::GPUModuleOp mlirModule, + const llvm::DataLayout &dl); /// Create an mlir::DataLayoutSpecInterface from the llvm.data_layout attribute /// if one is provided. If such attribute is not available, create a default @@ -37,6 +43,8 @@ void setMLIRDataLayout(mlir::ModuleOp mlirModule, const llvm::DataLayout &dl); /// nothing. void setMLIRDataLayoutFromAttributes(mlir::ModuleOp mlirModule, bool allowDefaultLayout); +void setMLIRDataLayoutFromAttributes(mlir::gpu::GPUModuleOp mlirModule, + bool allowDefaultLayout); /// Create mlir::DataLayout from the data layout information on the /// mlir::Module. Creates the data layout information attributes with @@ -44,7 +52,12 @@ void setMLIRDataLayoutFromAttributes(mlir::ModuleOp mlirModule, /// information is present at all and \p allowDefaultLayout is false, returns /// std::nullopt. std::optional -getOrSetDataLayout(mlir::ModuleOp mlirModule, bool allowDefaultLayout = false); +getOrSetMLIRDataLayout(mlir::ModuleOp mlirModule, + bool allowDefaultLayout = false); +std::optional +getOrSetMLIRDataLayout(mlir::gpu::GPUModuleOp mlirModule, + bool allowDefaultLayout = false); + } // namespace fir::support #endif // FORTRAN_OPTIMIZER_SUPPORT_DATALAYOUT_H diff --git a/flang/lib/Optimizer/CodeGen/CodeGen.cpp b/flang/lib/Optimizer/CodeGen/CodeGen.cpp index 1b078be7bb1c..cb4eb8303a49 100644 --- a/flang/lib/Optimizer/CodeGen/CodeGen.cpp +++ b/flang/lib/Optimizer/CodeGen/CodeGen.cpp @@ -1190,7 +1190,7 @@ genCUFAllocDescriptor(mlir::Location loc, mlir::ModuleOp mod, fir::BaseBoxType boxTy, const fir::LLVMTypeConverter &typeConverter) { std::optional dl = - fir::support::getOrSetDataLayout(mod, /*allowDefaultLayout=*/true); + fir::support::getOrSetMLIRDataLayout(mod, /*allowDefaultLayout=*/true); if (!dl) mlir::emitError(mod.getLoc(), "module operation must carry a data layout attribute " @@ -3942,7 +3942,7 @@ public: return signalPassFailure(); std::optional dl = - fir::support::getOrSetDataLayout(mod, /*allowDefaultLayout=*/true); + fir::support::getOrSetMLIRDataLayout(mod, /*allowDefaultLayout=*/true); if (!dl) { mlir::emitError(mod.getLoc(), "module operation must carry a data layout attribute " diff --git a/flang/lib/Optimizer/CodeGen/TargetRewrite.cpp b/flang/lib/Optimizer/CodeGen/TargetRewrite.cpp index b0b9499557e2..cc6c2b7df825 100644 --- a/flang/lib/Optimizer/CodeGen/TargetRewrite.cpp +++ b/flang/lib/Optimizer/CodeGen/TargetRewrite.cpp @@ -107,7 +107,7 @@ public: // TargetRewrite will require querying the type storage sizes, if it was // not set already, create a DataLayoutSpec for the ModuleOp now. std::optional dl = - fir::support::getOrSetDataLayout(mod, /*allowDefaultLayout=*/true); + fir::support::getOrSetMLIRDataLayout(mod, /*allowDefaultLayout=*/true); if (!dl) { mlir::emitError(mod.getLoc(), "module operation must carry a data layout attribute " diff --git a/flang/lib/Optimizer/Support/DataLayout.cpp b/flang/lib/Optimizer/Support/DataLayout.cpp index 93a3b92d0810..f89ce5984b91 100644 --- a/flang/lib/Optimizer/Support/DataLayout.cpp +++ b/flang/lib/Optimizer/Support/DataLayout.cpp @@ -10,6 +10,7 @@ #include "flang/Optimizer/Dialect/Support/FIRContext.h" #include "flang/Optimizer/Support/FatalError.h" #include "mlir/Dialect/DLTI/DLTI.h" +#include "mlir/Dialect/GPU/IR/GPUDialect.h" #include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/IR/BuiltinOps.h" #include "mlir/Interfaces/DataLayoutInterfaces.h" @@ -20,8 +21,9 @@ #include "llvm/Support/TargetSelect.h" #include "llvm/Target/TargetMachine.h" -void fir::support::setMLIRDataLayout(mlir::ModuleOp mlirModule, - const llvm::DataLayout &dl) { +namespace { +template +static void setDataLayout(ModOpTy mlirModule, const llvm::DataLayout &dl) { mlir::MLIRContext *context = mlirModule.getContext(); mlirModule->setAttr( mlir::LLVM::LLVMDialect::getDataLayoutAttrName(), @@ -30,12 +32,14 @@ void fir::support::setMLIRDataLayout(mlir::ModuleOp mlirModule, mlirModule->setAttr(mlir::DLTIDialect::kDataLayoutAttrName, dlSpec); } -void fir::support::setMLIRDataLayoutFromAttributes(mlir::ModuleOp mlirModule, - bool allowDefaultLayout) { +template +static void setDataLayoutFromAttributes(ModOpTy mlirModule, + bool allowDefaultLayout) { if (mlirModule.getDataLayoutSpec()) return; // Already set. - if (auto dataLayoutString = mlirModule->getAttrOfType( - mlir::LLVM::LLVMDialect::getDataLayoutAttrName())) { + if (auto dataLayoutString = + mlirModule->template getAttrOfType( + mlir::LLVM::LLVMDialect::getDataLayoutAttrName())) { llvm::DataLayout llvmDataLayout(dataLayoutString); fir::support::setMLIRDataLayout(mlirModule, llvmDataLayout); return; @@ -46,15 +50,48 @@ void fir::support::setMLIRDataLayoutFromAttributes(mlir::ModuleOp mlirModule, fir::support::setMLIRDataLayout(mlirModule, llvmDataLayout); } -std::optional -fir::support::getOrSetDataLayout(mlir::ModuleOp mlirModule, - bool allowDefaultLayout) { - if (!mlirModule.getDataLayoutSpec()) { +template +static std::optional +getOrSetDataLayout(ModOpTy mlirModule, bool allowDefaultLayout) { + if (!mlirModule.getDataLayoutSpec()) fir::support::setMLIRDataLayoutFromAttributes(mlirModule, allowDefaultLayout); - if (!mlirModule.getDataLayoutSpec()) { - return std::nullopt; - } - } + if (!mlirModule.getDataLayoutSpec() && + !mlir::isa(mlirModule)) + return std::nullopt; return mlir::DataLayout(mlirModule); } + +} // namespace + +void fir::support::setMLIRDataLayout(mlir::ModuleOp mlirModule, + const llvm::DataLayout &dl) { + setDataLayout(mlirModule, dl); +} + +void fir::support::setMLIRDataLayout(mlir::gpu::GPUModuleOp mlirModule, + const llvm::DataLayout &dl) { + setDataLayout(mlirModule, dl); +} + +void fir::support::setMLIRDataLayoutFromAttributes(mlir::ModuleOp mlirModule, + bool allowDefaultLayout) { + setDataLayoutFromAttributes(mlirModule, allowDefaultLayout); +} + +void fir::support::setMLIRDataLayoutFromAttributes( + mlir::gpu::GPUModuleOp mlirModule, bool allowDefaultLayout) { + setDataLayoutFromAttributes(mlirModule, allowDefaultLayout); +} + +std::optional +fir::support::getOrSetMLIRDataLayout(mlir::ModuleOp mlirModule, + bool allowDefaultLayout) { + return getOrSetDataLayout(mlirModule, allowDefaultLayout); +} + +std::optional +fir::support::getOrSetMLIRDataLayout(mlir::gpu::GPUModuleOp mlirModule, + bool allowDefaultLayout) { + return getOrSetDataLayout(mlirModule, allowDefaultLayout); +} diff --git a/flang/lib/Optimizer/Transforms/AddDebugInfo.cpp b/flang/lib/Optimizer/Transforms/AddDebugInfo.cpp index 16404fcda57b..3e1209c9ccd5 100644 --- a/flang/lib/Optimizer/Transforms/AddDebugInfo.cpp +++ b/flang/lib/Optimizer/Transforms/AddDebugInfo.cpp @@ -523,7 +523,7 @@ void AddDebugInfoPass::runOnOperation() { llvm::StringRef fileName; std::string filePath; std::optional dl = - fir::support::getOrSetDataLayout(module, /*allowDefaultLayout=*/true); + fir::support::getOrSetMLIRDataLayout(module, /*allowDefaultLayout=*/true); if (!dl) { mlir::emitError(module.getLoc(), "Missing data layout attribute in module"); signalPassFailure(); diff --git a/flang/lib/Optimizer/Transforms/CUFAddConstructor.cpp b/flang/lib/Optimizer/Transforms/CUFAddConstructor.cpp index 97551595db03..43ef6822de30 100644 --- a/flang/lib/Optimizer/Transforms/CUFAddConstructor.cpp +++ b/flang/lib/Optimizer/Transforms/CUFAddConstructor.cpp @@ -57,7 +57,7 @@ struct CUFAddConstructor auto funcTy = mlir::LLVM::LLVMFunctionType::get(voidTy, {}, /*isVarArg=*/false); std::optional dl = - fir::support::getOrSetDataLayout(mod, /*allowDefaultLayout=*/false); + fir::support::getOrSetMLIRDataLayout(mod, /*allowDefaultLayout=*/false); if (!dl) { mlir::emitError(mod.getLoc(), "data layout attribute is required to perform " + diff --git a/flang/lib/Optimizer/Transforms/CUFGPUToLLVMConversion.cpp b/flang/lib/Optimizer/Transforms/CUFGPUToLLVMConversion.cpp index c469b5a95b04..4611a18a541b 100644 --- a/flang/lib/Optimizer/Transforms/CUFGPUToLLVMConversion.cpp +++ b/flang/lib/Optimizer/Transforms/CUFGPUToLLVMConversion.cpp @@ -188,8 +188,8 @@ public: if (!module) return signalPassFailure(); - std::optional dl = - fir::support::getOrSetDataLayout(module, /*allowDefaultLayout=*/false); + std::optional dl = fir::support::getOrSetMLIRDataLayout( + module, /*allowDefaultLayout=*/false); fir::LLVMTypeConverter typeConverter(module, /*applyTBAA=*/false, /*forceUnifiedTBAATree=*/false, *dl); cuf::populateCUFGPUToLLVMConversionPatterns(typeConverter, patterns); diff --git a/flang/lib/Optimizer/Transforms/CUFOpConversion.cpp b/flang/lib/Optimizer/Transforms/CUFOpConversion.cpp index 710aed5031f5..77aa11f0603f 100644 --- a/flang/lib/Optimizer/Transforms/CUFOpConversion.cpp +++ b/flang/lib/Optimizer/Transforms/CUFOpConversion.cpp @@ -919,8 +919,8 @@ public: return signalPassFailure(); mlir::SymbolTable symtab(module); - std::optional dl = - fir::support::getOrSetDataLayout(module, /*allowDefaultLayout=*/false); + std::optional dl = fir::support::getOrSetMLIRDataLayout( + module, /*allowDefaultLayout=*/false); fir::LLVMTypeConverter typeConverter(module, /*applyTBAA=*/false, /*forceUnifiedTBAATree=*/false, *dl); target.addLegalDialectgetParentOfType(); fir::KindMapping kindMap = fir::getKindMapping(module); mlir::SmallVector argsOfInterest; - std::optional dl = - fir::support::getOrSetDataLayout(module, /*allowDefaultLayout=*/false); + std::optional dl = fir::support::getOrSetMLIRDataLayout( + module, /*allowDefaultLayout=*/false); if (!dl) mlir::emitError(module.getLoc(), "data layout attribute is required to perform " DEBUG_TYPE diff --git a/flang/test/lib/OpenACC/TestOpenACCInterfaces.cpp b/flang/test/lib/OpenACC/TestOpenACCInterfaces.cpp index a01396748f4c..5c14809a265e 100644 --- a/flang/test/lib/OpenACC/TestOpenACCInterfaces.cpp +++ b/flang/test/lib/OpenACC/TestOpenACCInterfaces.cpp @@ -28,7 +28,7 @@ struct TestFIROpenACCInterfaces void runOnOperation() override { mlir::ModuleOp mod = getOperation(); auto datalayout = - fir::support::getOrSetDataLayout(mod, /*allowDefaultLayout=*/true); + fir::support::getOrSetMLIRDataLayout(mod, /*allowDefaultLayout=*/true); mlir::OpBuilder builder(mod); getOperation().walk([&](Operation *op) { if (isa(op)) {