1 | //===-- Optimizer/Support/DataLayout.cpp ----------------------------------===// |
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/Support/DataLayout.h" |
10 | #include "flang/Optimizer/Dialect/Support/FIRContext.h" |
11 | #include "flang/Optimizer/Support/FatalError.h" |
12 | #include "mlir/Dialect/DLTI/DLTI.h" |
13 | #include "mlir/Dialect/GPU/IR/GPUDialect.h" |
14 | #include "mlir/Dialect/LLVMIR/LLVMDialect.h" |
15 | #include "mlir/IR/BuiltinOps.h" |
16 | #include "mlir/Interfaces/DataLayoutInterfaces.h" |
17 | #include "mlir/Support/LLVM.h" |
18 | #include "mlir/Target/LLVMIR/Import.h" |
19 | #include "llvm/IR/DataLayout.h" |
20 | #include "llvm/MC/TargetRegistry.h" |
21 | #include "llvm/Support/TargetSelect.h" |
22 | #include "llvm/Target/TargetMachine.h" |
23 | |
24 | namespace { |
25 | template <typename ModOpTy> |
26 | static void setDataLayout(ModOpTy mlirModule, const llvm::DataLayout &dl) { |
27 | mlir::MLIRContext *context = mlirModule.getContext(); |
28 | mlirModule->setAttr( |
29 | mlir::LLVM::LLVMDialect::getDataLayoutAttrName(), |
30 | mlir::StringAttr::get(context, dl.getStringRepresentation())); |
31 | mlir::DataLayoutSpecInterface dlSpec = mlir::translateDataLayout(dl, context); |
32 | mlirModule->setAttr(mlir::DLTIDialect::kDataLayoutAttrName, dlSpec); |
33 | } |
34 | |
35 | template <typename ModOpTy> |
36 | static void setDataLayoutFromAttributes(ModOpTy mlirModule, |
37 | bool allowDefaultLayout) { |
38 | if (mlirModule.getDataLayoutSpec()) |
39 | return; // Already set. |
40 | if (auto dataLayoutString = |
41 | mlirModule->template getAttrOfType<mlir::StringAttr>( |
42 | mlir::LLVM::LLVMDialect::getDataLayoutAttrName())) { |
43 | llvm::DataLayout llvmDataLayout(dataLayoutString); |
44 | fir::support::setMLIRDataLayout(mlirModule, llvmDataLayout); |
45 | return; |
46 | } |
47 | if (!allowDefaultLayout) |
48 | return; |
49 | llvm::DataLayout llvmDataLayout("" ); |
50 | fir::support::setMLIRDataLayout(mlirModule, llvmDataLayout); |
51 | } |
52 | |
53 | template <typename ModOpTy> |
54 | static std::optional<mlir::DataLayout> |
55 | getOrSetDataLayout(ModOpTy mlirModule, bool allowDefaultLayout) { |
56 | if (!mlirModule.getDataLayoutSpec()) |
57 | fir::support::setMLIRDataLayoutFromAttributes(mlirModule, |
58 | allowDefaultLayout); |
59 | if (!mlirModule.getDataLayoutSpec() && |
60 | !mlir::isa<mlir::gpu::GPUModuleOp>(mlirModule)) |
61 | return std::nullopt; |
62 | return mlir::DataLayout(mlirModule); |
63 | } |
64 | |
65 | } // namespace |
66 | |
67 | void fir::support::setMLIRDataLayout(mlir::ModuleOp mlirModule, |
68 | const llvm::DataLayout &dl) { |
69 | setDataLayout(mlirModule, dl); |
70 | } |
71 | |
72 | void fir::support::setMLIRDataLayout(mlir::gpu::GPUModuleOp mlirModule, |
73 | const llvm::DataLayout &dl) { |
74 | setDataLayout(mlirModule, dl); |
75 | } |
76 | |
77 | void fir::support::setMLIRDataLayoutFromAttributes(mlir::ModuleOp mlirModule, |
78 | bool allowDefaultLayout) { |
79 | setDataLayoutFromAttributes(mlirModule, allowDefaultLayout); |
80 | } |
81 | |
82 | void fir::support::setMLIRDataLayoutFromAttributes( |
83 | mlir::gpu::GPUModuleOp mlirModule, bool allowDefaultLayout) { |
84 | setDataLayoutFromAttributes(mlirModule, allowDefaultLayout); |
85 | } |
86 | |
87 | std::optional<mlir::DataLayout> |
88 | fir::support::getOrSetMLIRDataLayout(mlir::ModuleOp mlirModule, |
89 | bool allowDefaultLayout) { |
90 | return getOrSetDataLayout(mlirModule, allowDefaultLayout); |
91 | } |
92 | |
93 | std::optional<mlir::DataLayout> |
94 | fir::support::getOrSetMLIRDataLayout(mlir::gpu::GPUModuleOp mlirModule, |
95 | bool allowDefaultLayout) { |
96 | return getOrSetDataLayout(mlirModule, allowDefaultLayout); |
97 | } |
98 | |