[flang-commits] [flang] [mlir] [mlir][NFC] Mark type converter in `populate...` functions as `const` (PR #111250)
Matthias Springer via flang-commits
flang-commits at lists.llvm.org
Sat Oct 5 05:03:24 PDT 2024
https://github.com/matthias-springer created https://github.com/llvm/llvm-project/pull/111250
This commit marks the type converter in `populate...` functions as `const`. This is useful for debugging.
Patterns already take a `const` type converter. However, some `populate...` functions do not only add new patterns, but also add additional type conversion rules. That makes it difficult to find the place where a type conversion was added in the code base. With this change, all `populate...` functions that only populate pattern now have a `const` type converter. Programmers can then conclude from the function signature that these functions do not register any new type conversion rules.
Also some minor cleanups around the 1:N dialect conversion infrastructure, which did not always pass the type converter as a `const` object internally.
>From f69065c14b133bbfb38162c6e6094c10c001c157 Mon Sep 17 00:00:00 2001
From: Matthias Springer <mspringer at nvidia.com>
Date: Sat, 5 Oct 2024 13:56:58 +0200
Subject: [PATCH] [mlir][NFC] Mark type converter in `populate...` functions as
`const`
This commit marks the type converter in `populate...` functions as `const`. This is useful for debugging. Patterns already take a `const` type converter, but some `populate...` patterns do not only add new patterns, but also add additional type conversion rules. That makes it difficult to find the place where a type conversion was added in the code base. With this change, all `populate...` functions that only populate pattern now have a `const` type converter, from which programmers can conclude that they do not register any new type conversion rules.
Also some minor cleanups around the 1:N dialect conversion infrastructure, which did not always pass the type converter as a `const` object internally.
---
.../include/flang/Optimizer/CodeGen/CodeGen.h | 6 +-
.../flang/Optimizer/CodeGen/CodeGenOpenMP.h | 2 +-
.../Optimizer/Transforms/CufOpConversion.h | 2 +-
flang/lib/Optimizer/CodeGen/CodeGen.cpp | 2 +-
flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp | 2 +-
.../Optimizer/Transforms/CufOpConversion.cpp | 2 +-
mlir/docs/Bufferization.md | 2 +-
.../Conversion/AMDGPUToROCDL/AMDGPUToROCDL.h | 2 +-
.../mlir/Conversion/ArithToLLVM/ArithToLLVM.h | 2 +-
.../Conversion/ArithToSPIRV/ArithToSPIRV.h | 2 +-
.../Conversion/ComplexToLLVM/ComplexToLLVM.h | 2 +-
.../ComplexToSPIRV/ComplexToSPIRV.h | 2 +-
.../ControlFlowToLLVM/ControlFlowToLLVM.h | 6 +-
.../ControlFlowToSPIRV/ControlFlowToSPIRV.h | 2 +-
.../Conversion/FuncToLLVM/ConvertFuncToLLVM.h | 6 +-
.../mlir/Conversion/FuncToSPIRV/FuncToSPIRV.h | 2 +-
.../GPUToLLVMSPV/GPUToLLVMSPVPass.h | 2 +-
.../mlir/Conversion/GPUToNVVM/GPUToNVVMPass.h | 8 +-
.../Conversion/GPUToROCDL/GPUToROCDLPass.h | 2 +-
.../mlir/Conversion/GPUToSPIRV/GPUToSPIRV.h | 4 +-
.../mlir/Conversion/IndexToLLVM/IndexToLLVM.h | 2 +-
.../Conversion/IndexToSPIRV/IndexToSPIRV.h | 2 +-
.../mlir/Conversion/MathToLLVM/MathToLLVM.h | 2 +-
.../mlir/Conversion/MathToROCDL/MathToROCDL.h | 2 +-
.../mlir/Conversion/MathToSPIRV/MathToSPIRV.h | 2 +-
.../Conversion/MemRefToEmitC/MemRefToEmitC.h | 2 +-
.../Conversion/MemRefToLLVM/MemRefToLLVM.h | 2 +-
.../Conversion/MemRefToSPIRV/MemRefToSPIRV.h | 2 +-
.../mlir/Conversion/NVGPUToNVVM/NVGPUToNVVM.h | 2 +-
.../OpenMPToLLVM/ConvertOpenMPToLLVM.h | 4 +-
.../mlir/Conversion/SCFToSPIRV/SCFToSPIRV.h | 2 +-
.../mlir/Conversion/SPIRVToLLVM/SPIRVToLLVM.h | 17 +--
.../Conversion/TensorToSPIRV/TensorToSPIRV.h | 2 +-
.../Conversion/TosaToLinalg/TosaToLinalg.h | 2 +-
.../Conversion/TosaToTensor/TosaToTensor.h | 2 +-
.../mlir/Conversion/UBToLLVM/UBToLLVM.h | 2 +-
.../mlir/Conversion/UBToSPIRV/UBToSPIRV.h | 2 +-
.../VectorToLLVM/ConvertVectorToLLVM.h | 6 +-
.../Conversion/VectorToSPIRV/VectorToSPIRV.h | 2 +-
mlir/include/mlir/Dialect/AMX/Transforms.h | 4 +-
.../mlir/Dialect/Arith/Transforms/Passes.h | 10 +-
.../Dialect/ArmSVE/Transforms/Transforms.h | 4 +-
.../Bufferization/Transforms/Bufferize.h | 2 +-
.../Func/Transforms/DecomposeCallGraphTypes.h | 2 +-
.../Dialect/Func/Transforms/FuncConversions.h | 12 +-
.../Func/Transforms/OneToNFuncConversions.h | 2 +-
.../mlir/Dialect/Math/Transforms/Passes.h | 2 +-
.../Dialect/MemRef/Transforms/Transforms.h | 4 +-
.../mlir/Dialect/SCF/Transforms/Patterns.h | 8 +-
.../SPIRV/Transforms/SPIRVConversion.h | 2 +-
.../Dialect/SparseTensor/Transforms/Passes.h | 8 +-
.../Vector/Transforms/VectorRewritePatterns.h | 9 +-
.../mlir/Dialect/X86Vector/Transforms.h | 2 +-
.../mlir/Transforms/OneToNTypeConversion.h | 21 ++--
.../AMDGPUToROCDL/AMDGPUToROCDL.cpp | 20 ++--
.../Conversion/ArithToLLVM/ArithToLLVM.cpp | 2 +-
.../Conversion/ArithToSPIRV/ArithToSPIRV.cpp | 2 +-
.../ComplexToLLVM/ComplexToLLVM.cpp | 2 +-
.../ComplexToSPIRV/ComplexToSPIRV.cpp | 4 +-
.../ControlFlowToLLVM/ControlFlowToLLVM.cpp | 6 +-
.../ControlFlowToSPIRV/ControlFlowToSPIRV.cpp | 2 +-
.../ConvertToSPIRV/ConvertToSPIRVPass.cpp | 2 +-
mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp | 4 +-
.../Conversion/FuncToSPIRV/FuncToSPIRV.cpp | 2 +-
.../GPUCommon/IndexIntrinsicsOpLowering.h | 4 +-
.../GPUCommon/OpToFuncCallLowering.h | 2 +-
.../Conversion/GPUToLLVMSPV/GPUToLLVMSPV.cpp | 4 +-
.../GPUToNVVM/LowerGpuOpsToNVVMOps.cpp | 8 +-
.../Conversion/GPUToNVVM/WmmaOpsToNvvm.cpp | 2 +-
.../GPUToROCDL/LowerGpuOpsToROCDLOps.cpp | 4 +-
mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp | 5 +-
.../Conversion/GPUToSPIRV/WmmaOpsToSPIRV.cpp | 2 +-
.../Conversion/IndexToLLVM/IndexToLLVM.cpp | 2 +-
.../Conversion/IndexToSPIRV/IndexToSPIRV.cpp | 4 +-
mlir/lib/Conversion/MathToLLVM/MathToLLVM.cpp | 6 +-
.../Conversion/MathToROCDL/MathToROCDL.cpp | 6 +-
.../Conversion/MathToSPIRV/MathToSPIRV.cpp | 2 +-
.../MemRefToEmitC/MemRefToEmitC.cpp | 4 +-
.../Conversion/MemRefToLLVM/MemRefToLLVM.cpp | 2 +-
.../MemRefToSPIRV/MemRefToSPIRV.cpp | 2 +-
.../Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp | 4 +-
.../Conversion/OpenMPToLLVM/OpenMPToLLVM.cpp | 2 +-
mlir/lib/Conversion/SCFToSPIRV/SCFToSPIRV.cpp | 6 +-
.../Conversion/SPIRVToLLVM/SPIRVToLLVM.cpp | 107 +++++++++---------
.../TensorToSPIRV/TensorToSPIRV.cpp | 8 +-
.../Conversion/TosaToLinalg/TosaToLinalg.cpp | 2 +-
.../Conversion/TosaToTensor/TosaToTensor.cpp | 73 ++++++------
mlir/lib/Conversion/UBToLLVM/UBToLLVM.cpp | 4 +-
mlir/lib/Conversion/UBToSPIRV/UBToSPIRV.cpp | 2 +-
.../VectorToLLVM/ConvertVectorToLLVM.cpp | 4 +-
.../VectorToSPIRV/VectorToSPIRV.cpp | 4 +-
.../AMX/Transforms/LegalizeForLLVMExport.cpp | 2 +-
.../Arith/Transforms/EmulateNarrowType.cpp | 3 +-
.../Transforms/EmulateUnsupportedFloats.cpp | 6 +-
.../Arith/Transforms/EmulateWideInt.cpp | 3 +-
.../Transforms/LegalizeForLLVMExport.cpp | 2 +-
.../Bufferization/Transforms/Bufferize.cpp | 2 +-
.../Transforms/DecomposeCallGraphTypes.cpp | 4 +-
.../Func/Transforms/FuncConversions.cpp | 17 ++-
.../Func/Transforms/OneToNFuncConversions.cpp | 2 +-
.../Transforms/ExtendToSupportedTypes.cpp | 4 +-
.../MemRef/Transforms/EmulateNarrowType.cpp | 2 +-
.../MemRef/Transforms/EmulateWideInt.cpp | 2 +-
.../SCF/Transforms/OneToNTypeConversion.cpp | 4 +-
.../Transforms/StructuralTypeConversions.cpp | 4 +-
.../SPIRV/Transforms/SPIRVConversion.cpp | 4 +-
.../Transforms/SparseIterationToScf.cpp | 2 +-
.../SparseStorageSpecifierToLLVM.cpp | 4 +-
.../Transforms/SparseTensorCodegen.cpp | 14 +--
.../Transforms/SparseTensorConversion.cpp | 4 +-
.../Transforms/VectorEmulateNarrowType.cpp | 2 +-
.../Vector/Transforms/VectorLinearize.cpp | 2 +-
.../Transforms/LegalizeForLLVMExport.cpp | 6 +-
.../Transforms/Utils/OneToNTypeConversion.cpp | 4 +-
.../TestOneToNTypeConversionPass.cpp | 5 +-
115 files changed, 327 insertions(+), 309 deletions(-)
diff --git a/flang/include/flang/Optimizer/CodeGen/CodeGen.h b/flang/include/flang/Optimizer/CodeGen/CodeGen.h
index 390f00e1ac77c2..255b1950c84256 100644
--- a/flang/include/flang/Optimizer/CodeGen/CodeGen.h
+++ b/flang/include/flang/Optimizer/CodeGen/CodeGen.h
@@ -72,9 +72,9 @@ std::unique_ptr<mlir::Pass> createLLVMDialectToLLVMPass(
[](llvm::Module &m, llvm::raw_ostream &out) { m.print(out, nullptr); });
/// Populate the given list with patterns that convert from FIR to LLVM.
-void populateFIRToLLVMConversionPatterns(fir::LLVMTypeConverter &converter,
- mlir::RewritePatternSet &patterns,
- fir::FIRToLLVMPassOptions &options);
+void populateFIRToLLVMConversionPatterns(
+ const fir::LLVMTypeConverter &converter, mlir::RewritePatternSet &patterns,
+ fir::FIRToLLVMPassOptions &options);
/// Populate the pattern set with the PreCGRewrite patterns.
void populatePreCGRewritePatterns(mlir::RewritePatternSet &patterns,
diff --git a/flang/include/flang/Optimizer/CodeGen/CodeGenOpenMP.h b/flang/include/flang/Optimizer/CodeGen/CodeGenOpenMP.h
index 1832d4967b7ecd..04b59da72f6bae 100644
--- a/flang/include/flang/Optimizer/CodeGen/CodeGenOpenMP.h
+++ b/flang/include/flang/Optimizer/CodeGen/CodeGenOpenMP.h
@@ -19,7 +19,7 @@ class LLVMTypeConverter;
/// dialect, utilised in cases where the default OpenMP dialect handling cannot
/// handle all cases for intermingled fir types and operations.
void populateOpenMPFIRToLLVMConversionPatterns(
- LLVMTypeConverter &converter, mlir::RewritePatternSet &patterns);
+ const LLVMTypeConverter &converter, mlir::RewritePatternSet &patterns);
} // namespace fir
diff --git a/flang/include/flang/Optimizer/Transforms/CufOpConversion.h b/flang/include/flang/Optimizer/Transforms/CufOpConversion.h
index 3e13c72cc0e3e4..79ce4ac5c6cbc0 100644
--- a/flang/include/flang/Optimizer/Transforms/CufOpConversion.h
+++ b/flang/include/flang/Optimizer/Transforms/CufOpConversion.h
@@ -22,7 +22,7 @@ class DataLayout;
namespace cuf {
-void populateCUFToFIRConversionPatterns(fir::LLVMTypeConverter &converter,
+void populateCUFToFIRConversionPatterns(const fir::LLVMTypeConverter &converter,
mlir::DataLayout &dl,
mlir::RewritePatternSet &patterns);
diff --git a/flang/lib/Optimizer/CodeGen/CodeGen.cpp b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
index 1cb869bfeb95a8..1611de9e6389a5 100644
--- a/flang/lib/Optimizer/CodeGen/CodeGen.cpp
+++ b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
@@ -3823,7 +3823,7 @@ fir::createLLVMDialectToLLVMPass(llvm::raw_ostream &output,
}
void fir::populateFIRToLLVMConversionPatterns(
- fir::LLVMTypeConverter &converter, mlir::RewritePatternSet &patterns,
+ const fir::LLVMTypeConverter &converter, mlir::RewritePatternSet &patterns,
fir::FIRToLLVMPassOptions &options) {
patterns.insert<
AbsentOpConversion, AddcOpConversion, AddrOfOpConversion,
diff --git a/flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp b/flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp
index 48d6c000ee44a3..da13ed648e44e6 100644
--- a/flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp
+++ b/flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp
@@ -93,6 +93,6 @@ struct MapInfoOpConversion
} // namespace
void fir::populateOpenMPFIRToLLVMConversionPatterns(
- LLVMTypeConverter &converter, mlir::RewritePatternSet &patterns) {
+ const LLVMTypeConverter &converter, mlir::RewritePatternSet &patterns) {
patterns.add<MapInfoOpConversion>(converter);
}
diff --git a/flang/lib/Optimizer/Transforms/CufOpConversion.cpp b/flang/lib/Optimizer/Transforms/CufOpConversion.cpp
index 24dc8e2354b120..c573506e233599 100644
--- a/flang/lib/Optimizer/Transforms/CufOpConversion.cpp
+++ b/flang/lib/Optimizer/Transforms/CufOpConversion.cpp
@@ -583,7 +583,7 @@ class CufOpConversion : public fir::impl::CufOpConversionBase<CufOpConversion> {
} // namespace
void cuf::populateCUFToFIRConversionPatterns(
- fir::LLVMTypeConverter &converter, mlir::DataLayout &dl,
+ const fir::LLVMTypeConverter &converter, mlir::DataLayout &dl,
mlir::RewritePatternSet &patterns) {
patterns.insert<CufAllocOpConversion>(patterns.getContext(), &dl, &converter);
patterns.insert<CufAllocateOpConversion, CufDeallocateOpConversion,
diff --git a/mlir/docs/Bufferization.md b/mlir/docs/Bufferization.md
index 6a49bea9c68c5d..d5a426e09e7ceb 100644
--- a/mlir/docs/Bufferization.md
+++ b/mlir/docs/Bufferization.md
@@ -651,7 +651,7 @@ is very small, and follows the basic pattern of any dialect conversion pass.
```
void mlir::populateTensorBufferizePatterns(
- BufferizeTypeConverter &typeConverter, RewritePatternSet &patterns) {
+ const BufferizeTypeConverter &typeConverter, RewritePatternSet &patterns) {
patterns.add<BufferizeCastOp, BufferizeExtractOp>(typeConverter,
patterns.getContext());
}
diff --git a/mlir/include/mlir/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.h b/mlir/include/mlir/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.h
index e445c58fe6c547..e7637a6013e68a 100644
--- a/mlir/include/mlir/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.h
+++ b/mlir/include/mlir/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.h
@@ -24,7 +24,7 @@ class Pass;
/// Note: The ROCDL target does not support the LLVM bfloat type at this time
/// and so this function will add conversions to change all `bfloat` uses
/// to `i16`.
-void populateAMDGPUToROCDLConversionPatterns(LLVMTypeConverter &converter,
+void populateAMDGPUToROCDLConversionPatterns(const LLVMTypeConverter &converter,
RewritePatternSet &patterns,
amdgpu::Chipset chipset);
diff --git a/mlir/include/mlir/Conversion/ArithToLLVM/ArithToLLVM.h b/mlir/include/mlir/Conversion/ArithToLLVM/ArithToLLVM.h
index 25b82aed9c968a..173fea7c352d93 100644
--- a/mlir/include/mlir/Conversion/ArithToLLVM/ArithToLLVM.h
+++ b/mlir/include/mlir/Conversion/ArithToLLVM/ArithToLLVM.h
@@ -22,7 +22,7 @@ class Pass;
#include "mlir/Conversion/Passes.h.inc"
namespace arith {
-void populateArithToLLVMConversionPatterns(LLVMTypeConverter &converter,
+void populateArithToLLVMConversionPatterns(const LLVMTypeConverter &converter,
RewritePatternSet &patterns);
void registerConvertArithToLLVMInterface(DialectRegistry ®istry);
diff --git a/mlir/include/mlir/Conversion/ArithToSPIRV/ArithToSPIRV.h b/mlir/include/mlir/Conversion/ArithToSPIRV/ArithToSPIRV.h
index bb30deb9dc10e5..a15e56bfdf04bf 100644
--- a/mlir/include/mlir/Conversion/ArithToSPIRV/ArithToSPIRV.h
+++ b/mlir/include/mlir/Conversion/ArithToSPIRV/ArithToSPIRV.h
@@ -22,7 +22,7 @@ class Pass;
#include "mlir/Conversion/Passes.h.inc"
namespace arith {
-void populateArithToSPIRVPatterns(SPIRVTypeConverter &typeConverter,
+void populateArithToSPIRVPatterns(const SPIRVTypeConverter &typeConverter,
RewritePatternSet &patterns);
std::unique_ptr<OperationPass<>> createConvertArithToSPIRVPass();
diff --git a/mlir/include/mlir/Conversion/ComplexToLLVM/ComplexToLLVM.h b/mlir/include/mlir/Conversion/ComplexToLLVM/ComplexToLLVM.h
index 69b01b471dd984..65d6f165c7448d 100644
--- a/mlir/include/mlir/Conversion/ComplexToLLVM/ComplexToLLVM.h
+++ b/mlir/include/mlir/Conversion/ComplexToLLVM/ComplexToLLVM.h
@@ -39,7 +39,7 @@ class ComplexStructBuilder : public StructBuilder {
};
/// Populate the given list with patterns that convert from Complex to LLVM.
-void populateComplexToLLVMConversionPatterns(LLVMTypeConverter &converter,
+void populateComplexToLLVMConversionPatterns(const LLVMTypeConverter &converter,
RewritePatternSet &patterns);
void registerConvertComplexToLLVMInterface(DialectRegistry ®istry);
diff --git a/mlir/include/mlir/Conversion/ComplexToSPIRV/ComplexToSPIRV.h b/mlir/include/mlir/Conversion/ComplexToSPIRV/ComplexToSPIRV.h
index 19dba5b870cb6f..777481f00b3592 100644
--- a/mlir/include/mlir/Conversion/ComplexToSPIRV/ComplexToSPIRV.h
+++ b/mlir/include/mlir/Conversion/ComplexToSPIRV/ComplexToSPIRV.h
@@ -20,7 +20,7 @@ class SPIRVTypeConverter;
/// Appends to a pattern list additional patterns for translating Complex ops
/// to SPIR-V ops.
-void populateComplexToSPIRVPatterns(SPIRVTypeConverter &typeConverter,
+void populateComplexToSPIRVPatterns(const SPIRVTypeConverter &typeConverter,
RewritePatternSet &patterns);
} // namespace mlir
diff --git a/mlir/include/mlir/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.h b/mlir/include/mlir/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.h
index 35b691b8bc44b5..b88c1e8b20f32b 100644
--- a/mlir/include/mlir/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.h
+++ b/mlir/include/mlir/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.h
@@ -29,13 +29,13 @@ namespace cf {
/// Collect the patterns to convert from the ControlFlow dialect to LLVM. The
/// conversion patterns capture the LLVMTypeConverter by reference meaning the
/// references have to remain alive during the entire pattern lifetime.
-void populateControlFlowToLLVMConversionPatterns(LLVMTypeConverter &converter,
- RewritePatternSet &patterns);
+void populateControlFlowToLLVMConversionPatterns(
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns);
/// Populate the cf.assert to LLVM conversion pattern. If `abortOnFailure` is
/// set to false, the program execution continues when a condition is
/// unsatisfied.
-void populateAssertToLLVMConversionPattern(LLVMTypeConverter &converter,
+void populateAssertToLLVMConversionPattern(const LLVMTypeConverter &converter,
RewritePatternSet &patterns,
bool abortOnFailure = true);
diff --git a/mlir/include/mlir/Conversion/ControlFlowToSPIRV/ControlFlowToSPIRV.h b/mlir/include/mlir/Conversion/ControlFlowToSPIRV/ControlFlowToSPIRV.h
index 43578ffffae2dc..6de46960ff8029 100644
--- a/mlir/include/mlir/Conversion/ControlFlowToSPIRV/ControlFlowToSPIRV.h
+++ b/mlir/include/mlir/Conversion/ControlFlowToSPIRV/ControlFlowToSPIRV.h
@@ -20,7 +20,7 @@ class SPIRVTypeConverter;
namespace cf {
/// Appends to a pattern list additional patterns for translating ControlFLow
/// ops to SPIR-V ops.
-void populateControlFlowToSPIRVPatterns(SPIRVTypeConverter &typeConverter,
+void populateControlFlowToSPIRVPatterns(const SPIRVTypeConverter &typeConverter,
RewritePatternSet &patterns);
} // namespace cf
} // namespace mlir
diff --git a/mlir/include/mlir/Conversion/FuncToLLVM/ConvertFuncToLLVM.h b/mlir/include/mlir/Conversion/FuncToLLVM/ConvertFuncToLLVM.h
index 76eb3b5dd11e95..b1ea2740c06059 100644
--- a/mlir/include/mlir/Conversion/FuncToLLVM/ConvertFuncToLLVM.h
+++ b/mlir/include/mlir/Conversion/FuncToLLVM/ConvertFuncToLLVM.h
@@ -39,8 +39,8 @@ convertFuncOpToLLVMFuncOp(FunctionOpInterface funcOp,
/// `emitCWrappers` is set, the pattern will also produce functions
/// that pass memref descriptors by pointer-to-structure in addition to the
/// default unpacked form.
-void populateFuncToLLVMFuncOpConversionPattern(LLVMTypeConverter &converter,
- RewritePatternSet &patterns);
+void populateFuncToLLVMFuncOpConversionPattern(
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns);
/// Collect the patterns to convert from the Func dialect to LLVM. The
/// conversion patterns capture the LLVMTypeConverter and the LowerToLLVMOptions
@@ -56,7 +56,7 @@ void populateFuncToLLVMFuncOpConversionPattern(LLVMTypeConverter &converter,
/// needed if `converter.getOptions().useBarePtrCallConv` is `true`, but it's
/// not an error to provide it anyway.
void populateFuncToLLVMConversionPatterns(
- LLVMTypeConverter &converter, RewritePatternSet &patterns,
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns,
const SymbolTable *symbolTable = nullptr);
void registerConvertFuncToLLVMInterface(DialectRegistry ®istry);
diff --git a/mlir/include/mlir/Conversion/FuncToSPIRV/FuncToSPIRV.h b/mlir/include/mlir/Conversion/FuncToSPIRV/FuncToSPIRV.h
index 2fa55f40dd9700..68728755a2fbdb 100644
--- a/mlir/include/mlir/Conversion/FuncToSPIRV/FuncToSPIRV.h
+++ b/mlir/include/mlir/Conversion/FuncToSPIRV/FuncToSPIRV.h
@@ -21,7 +21,7 @@ class SPIRVTypeConverter;
/// Appends to a pattern list additional patterns for translating Func ops
/// to SPIR-V ops. Also adds the patterns to legalize ops not directly
/// translated to SPIR-V dialect.
-void populateFuncToSPIRVPatterns(SPIRVTypeConverter &typeConverter,
+void populateFuncToSPIRVPatterns(const SPIRVTypeConverter &typeConverter,
RewritePatternSet &patterns);
} // namespace mlir
diff --git a/mlir/include/mlir/Conversion/GPUToLLVMSPV/GPUToLLVMSPVPass.h b/mlir/include/mlir/Conversion/GPUToLLVMSPV/GPUToLLVMSPVPass.h
index 20c9f9c394525d..3bf30ae99f2bf1 100644
--- a/mlir/include/mlir/Conversion/GPUToLLVMSPV/GPUToLLVMSPVPass.h
+++ b/mlir/include/mlir/Conversion/GPUToLLVMSPV/GPUToLLVMSPVPass.h
@@ -21,7 +21,7 @@ class TypeConverter;
#define GEN_PASS_DECL_CONVERTGPUOPSTOLLVMSPVOPS
#include "mlir/Conversion/Passes.h.inc"
-void populateGpuToLLVMSPVConversionPatterns(LLVMTypeConverter &converter,
+void populateGpuToLLVMSPVConversionPatterns(const LLVMTypeConverter &converter,
RewritePatternSet &patterns);
/// Populates memory space attribute conversion rules for lowering
diff --git a/mlir/include/mlir/Conversion/GPUToNVVM/GPUToNVVMPass.h b/mlir/include/mlir/Conversion/GPUToNVVM/GPUToNVVMPass.h
index e0f4c71051e506..645e86a4309621 100644
--- a/mlir/include/mlir/Conversion/GPUToNVVM/GPUToNVVMPass.h
+++ b/mlir/include/mlir/Conversion/GPUToNVVM/GPUToNVVMPass.h
@@ -32,16 +32,16 @@ LLVM::LLVMStructType convertMMAToLLVMType(gpu::MMAMatrixType type);
void configureGpuToNVVMConversionLegality(ConversionTarget &target);
/// Collect a set of patterns to convert from the GPU dialect to NVVM.
-void populateGpuToNVVMConversionPatterns(LLVMTypeConverter &converter,
+void populateGpuToNVVMConversionPatterns(const LLVMTypeConverter &converter,
RewritePatternSet &patterns);
/// Populate GpuSubgroupReduce pattern to NVVM. It generates a specific nvvm
/// op that is not available on every GPU.
-void populateGpuSubgroupReduceOpLoweringPattern(LLVMTypeConverter &converter,
- RewritePatternSet &patterns);
+void populateGpuSubgroupReduceOpLoweringPattern(
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns);
/// Collect a set of patterns to convert WMMA ops from GPU dialect to NVVM.
-void populateGpuWMMAToNVVMConversionPatterns(LLVMTypeConverter &converter,
+void populateGpuWMMAToNVVMConversionPatterns(const LLVMTypeConverter &converter,
RewritePatternSet &patterns);
} // namespace mlir
diff --git a/mlir/include/mlir/Conversion/GPUToROCDL/GPUToROCDLPass.h b/mlir/include/mlir/Conversion/GPUToROCDL/GPUToROCDLPass.h
index 5647787712997b..1a917932a9a84a 100644
--- a/mlir/include/mlir/Conversion/GPUToROCDL/GPUToROCDLPass.h
+++ b/mlir/include/mlir/Conversion/GPUToROCDL/GPUToROCDLPass.h
@@ -30,7 +30,7 @@ class GPUModuleOp;
/// Collect a set of patterns to convert from the GPU dialect to ROCDL.
/// If `runtime` is Unknown, gpu.printf will not be lowered
/// The resulting pattern set should be run over a gpu.module op
-void populateGpuToROCDLConversionPatterns(LLVMTypeConverter &converter,
+void populateGpuToROCDLConversionPatterns(const LLVMTypeConverter &converter,
RewritePatternSet &patterns,
gpu::amd::Runtime runtime);
diff --git a/mlir/include/mlir/Conversion/GPUToSPIRV/GPUToSPIRV.h b/mlir/include/mlir/Conversion/GPUToSPIRV/GPUToSPIRV.h
index d34549432161db..825142d120587f 100644
--- a/mlir/include/mlir/Conversion/GPUToSPIRV/GPUToSPIRV.h
+++ b/mlir/include/mlir/Conversion/GPUToSPIRV/GPUToSPIRV.h
@@ -23,13 +23,13 @@ class SPIRVTypeConverter;
/// Appends to a pattern list additional patterns for translating GPU Ops to
/// SPIR-V ops. For a gpu.func to be converted, it should have a
/// spirv.entry_point_abi attribute.
-void populateGPUToSPIRVPatterns(SPIRVTypeConverter &typeConverter,
+void populateGPUToSPIRVPatterns(const SPIRVTypeConverter &typeConverter,
RewritePatternSet &patterns);
/// Collect a set of patterns to convert WMMA ops from GPU dialect to SPIRV,
/// using the KHR Cooperative Matrix extension.
void populateGpuWMMAToSPIRVCoopMatrixKHRConversionPatterns(
- SPIRVTypeConverter &typeConverter, RewritePatternSet &patterns);
+ const SPIRVTypeConverter &typeConverter, RewritePatternSet &patterns);
/// Adds `MMAMatrixType` conversions to SPIR-V cooperative matrix KHR type
/// conversion to the type converter.
diff --git a/mlir/include/mlir/Conversion/IndexToLLVM/IndexToLLVM.h b/mlir/include/mlir/Conversion/IndexToLLVM/IndexToLLVM.h
index df605adae669be..16d3d74b135588 100644
--- a/mlir/include/mlir/Conversion/IndexToLLVM/IndexToLLVM.h
+++ b/mlir/include/mlir/Conversion/IndexToLLVM/IndexToLLVM.h
@@ -21,7 +21,7 @@ class Pass;
#include "mlir/Conversion/Passes.h.inc"
namespace index {
-void populateIndexToLLVMConversionPatterns(LLVMTypeConverter &converter,
+void populateIndexToLLVMConversionPatterns(const LLVMTypeConverter &converter,
RewritePatternSet &patterns);
void registerConvertIndexToLLVMInterface(DialectRegistry ®istry);
diff --git a/mlir/include/mlir/Conversion/IndexToSPIRV/IndexToSPIRV.h b/mlir/include/mlir/Conversion/IndexToSPIRV/IndexToSPIRV.h
index 58a1c5246eef99..e5e7d3318a8717 100644
--- a/mlir/include/mlir/Conversion/IndexToSPIRV/IndexToSPIRV.h
+++ b/mlir/include/mlir/Conversion/IndexToSPIRV/IndexToSPIRV.h
@@ -21,7 +21,7 @@ class Pass;
#include "mlir/Conversion/Passes.h.inc"
namespace index {
-void populateIndexToSPIRVPatterns(SPIRVTypeConverter &converter,
+void populateIndexToSPIRVPatterns(const SPIRVTypeConverter &converter,
RewritePatternSet &patterns);
std::unique_ptr<OperationPass<>> createConvertIndexToSPIRVPass();
} // namespace index
diff --git a/mlir/include/mlir/Conversion/MathToLLVM/MathToLLVM.h b/mlir/include/mlir/Conversion/MathToLLVM/MathToLLVM.h
index 4ddc1637000fdb..93cd780bba4381 100644
--- a/mlir/include/mlir/Conversion/MathToLLVM/MathToLLVM.h
+++ b/mlir/include/mlir/Conversion/MathToLLVM/MathToLLVM.h
@@ -21,7 +21,7 @@ class Pass;
#define GEN_PASS_DECL_CONVERTMATHTOLLVMPASS
#include "mlir/Conversion/Passes.h.inc"
-void populateMathToLLVMConversionPatterns(LLVMTypeConverter &converter,
+void populateMathToLLVMConversionPatterns(const LLVMTypeConverter &converter,
RewritePatternSet &patterns,
bool approximateLog1p = true);
diff --git a/mlir/include/mlir/Conversion/MathToROCDL/MathToROCDL.h b/mlir/include/mlir/Conversion/MathToROCDL/MathToROCDL.h
index fa7a635568c7c2..46573e7966cccc 100644
--- a/mlir/include/mlir/Conversion/MathToROCDL/MathToROCDL.h
+++ b/mlir/include/mlir/Conversion/MathToROCDL/MathToROCDL.h
@@ -19,7 +19,7 @@ class Pass;
#include "mlir/Conversion/Passes.h.inc"
/// Populate the given list with patterns that convert from Math to ROCDL calls.
-void populateMathToROCDLConversionPatterns(LLVMTypeConverter &converter,
+void populateMathToROCDLConversionPatterns(const LLVMTypeConverter &converter,
RewritePatternSet &patterns);
} // namespace mlir
diff --git a/mlir/include/mlir/Conversion/MathToSPIRV/MathToSPIRV.h b/mlir/include/mlir/Conversion/MathToSPIRV/MathToSPIRV.h
index 10090268a46636..452b3f5df931e1 100644
--- a/mlir/include/mlir/Conversion/MathToSPIRV/MathToSPIRV.h
+++ b/mlir/include/mlir/Conversion/MathToSPIRV/MathToSPIRV.h
@@ -20,7 +20,7 @@ class SPIRVTypeConverter;
/// Appends to a pattern list additional patterns for translating Math ops
/// to SPIR-V ops.
-void populateMathToSPIRVPatterns(SPIRVTypeConverter &typeConverter,
+void populateMathToSPIRVPatterns(const SPIRVTypeConverter &typeConverter,
RewritePatternSet &patterns);
} // namespace mlir
diff --git a/mlir/include/mlir/Conversion/MemRefToEmitC/MemRefToEmitC.h b/mlir/include/mlir/Conversion/MemRefToEmitC/MemRefToEmitC.h
index 734ffdba520c9f..950d50229bac16 100644
--- a/mlir/include/mlir/Conversion/MemRefToEmitC/MemRefToEmitC.h
+++ b/mlir/include/mlir/Conversion/MemRefToEmitC/MemRefToEmitC.h
@@ -15,7 +15,7 @@ class TypeConverter;
void populateMemRefToEmitCTypeConversion(TypeConverter &typeConverter);
void populateMemRefToEmitCConversionPatterns(RewritePatternSet &patterns,
- TypeConverter &converter);
+ const TypeConverter &converter);
} // namespace mlir
#endif // MLIR_CONVERSION_MEMREFTOEMITC_MEMREFTOEMITC_H
diff --git a/mlir/include/mlir/Conversion/MemRefToLLVM/MemRefToLLVM.h b/mlir/include/mlir/Conversion/MemRefToLLVM/MemRefToLLVM.h
index db1c222643d571..996a64baf9dd58 100644
--- a/mlir/include/mlir/Conversion/MemRefToLLVM/MemRefToLLVM.h
+++ b/mlir/include/mlir/Conversion/MemRefToLLVM/MemRefToLLVM.h
@@ -23,7 +23,7 @@ class RewritePatternSet;
/// Collect a set of patterns to convert memory-related operations from the
/// MemRef dialect to the LLVM dialect.
void populateFinalizeMemRefToLLVMConversionPatterns(
- LLVMTypeConverter &converter, RewritePatternSet &patterns);
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns);
void registerConvertMemRefToLLVMInterface(DialectRegistry ®istry);
diff --git a/mlir/include/mlir/Conversion/MemRefToSPIRV/MemRefToSPIRV.h b/mlir/include/mlir/Conversion/MemRefToSPIRV/MemRefToSPIRV.h
index 54711c8ad727fb..4aa2bac98a7bb9 100644
--- a/mlir/include/mlir/Conversion/MemRefToSPIRV/MemRefToSPIRV.h
+++ b/mlir/include/mlir/Conversion/MemRefToSPIRV/MemRefToSPIRV.h
@@ -67,7 +67,7 @@ void convertMemRefTypesAndAttrs(
/// Appends to a pattern list additional patterns for translating MemRef ops
/// to SPIR-V ops.
-void populateMemRefToSPIRVPatterns(SPIRVTypeConverter &typeConverter,
+void populateMemRefToSPIRVPatterns(const SPIRVTypeConverter &typeConverter,
RewritePatternSet &patterns);
} // namespace mlir
diff --git a/mlir/include/mlir/Conversion/NVGPUToNVVM/NVGPUToNVVM.h b/mlir/include/mlir/Conversion/NVGPUToNVVM/NVGPUToNVVM.h
index 4b8d5c5fe2a893..ee5b8cecb529c4 100644
--- a/mlir/include/mlir/Conversion/NVGPUToNVVM/NVGPUToNVVM.h
+++ b/mlir/include/mlir/Conversion/NVGPUToNVVM/NVGPUToNVVM.h
@@ -34,7 +34,7 @@ MemRefType getMBarrierMemrefType(MLIRContext *context,
MBarrierGroupType barrierType);
} // namespace nvgpu
-void populateNVGPUToNVVMConversionPatterns(LLVMTypeConverter &converter,
+void populateNVGPUToNVVMConversionPatterns(const LLVMTypeConverter &converter,
RewritePatternSet &patterns);
} // namespace mlir
diff --git a/mlir/include/mlir/Conversion/OpenMPToLLVM/ConvertOpenMPToLLVM.h b/mlir/include/mlir/Conversion/OpenMPToLLVM/ConvertOpenMPToLLVM.h
index 2ed077f8f83e91..5ae484eadf49ff 100644
--- a/mlir/include/mlir/Conversion/OpenMPToLLVM/ConvertOpenMPToLLVM.h
+++ b/mlir/include/mlir/Conversion/OpenMPToLLVM/ConvertOpenMPToLLVM.h
@@ -22,8 +22,8 @@ class RewritePatternSet;
/// Configure dynamic conversion legality of regionless operations from OpenMP
/// to LLVM.
-void configureOpenMPToLLVMConversionLegality(ConversionTarget &target,
- LLVMTypeConverter &typeConverter);
+void configureOpenMPToLLVMConversionLegality(
+ ConversionTarget &target, const LLVMTypeConverter &typeConverter);
/// Populate the given list with patterns that convert from OpenMP to LLVM.
void populateOpenMPToLLVMConversionPatterns(LLVMTypeConverter &converter,
diff --git a/mlir/include/mlir/Conversion/SCFToSPIRV/SCFToSPIRV.h b/mlir/include/mlir/Conversion/SCFToSPIRV/SCFToSPIRV.h
index 4effa7186ef36f..5b960cde1050cb 100644
--- a/mlir/include/mlir/Conversion/SCFToSPIRV/SCFToSPIRV.h
+++ b/mlir/include/mlir/Conversion/SCFToSPIRV/SCFToSPIRV.h
@@ -34,7 +34,7 @@ struct ScfToSPIRVContext {
/// Collects a set of patterns to lower from scf.for, scf.if, and
/// loop.terminator to CFG operations within the SPIR-V dialect.
-void populateSCFToSPIRVPatterns(SPIRVTypeConverter &typeConverter,
+void populateSCFToSPIRVPatterns(const SPIRVTypeConverter &typeConverter,
ScfToSPIRVContext &scfToSPIRVContext,
RewritePatternSet &patterns);
} // namespace mlir
diff --git a/mlir/include/mlir/Conversion/SPIRVToLLVM/SPIRVToLLVM.h b/mlir/include/mlir/Conversion/SPIRVToLLVM/SPIRVToLLVM.h
index 84935d19670437..e59373c955fc3b 100644
--- a/mlir/include/mlir/Conversion/SPIRVToLLVM/SPIRVToLLVM.h
+++ b/mlir/include/mlir/Conversion/SPIRVToLLVM/SPIRVToLLVM.h
@@ -25,13 +25,16 @@ class ModuleOp;
template <typename SPIRVOp>
class SPIRVToLLVMConversion : public OpConversionPattern<SPIRVOp> {
public:
- SPIRVToLLVMConversion(MLIRContext *context, LLVMTypeConverter &typeConverter,
+ SPIRVToLLVMConversion(MLIRContext *context,
+ const LLVMTypeConverter &typeConverter,
PatternBenefit benefit = 1)
- : OpConversionPattern<SPIRVOp>(typeConverter, context, benefit),
- typeConverter(typeConverter) {}
+ : OpConversionPattern<SPIRVOp>(typeConverter, context, benefit) {}
protected:
- LLVMTypeConverter &typeConverter;
+ const LLVMTypeConverter *getTypeConverter() const {
+ return static_cast<const LLVMTypeConverter *>(
+ ConversionPattern::getTypeConverter());
+ }
};
/// Encodes global variable's descriptor set and binding into its name if they
@@ -46,18 +49,18 @@ void populateSPIRVToLLVMTypeConversion(
/// Populates the given list with patterns that convert from SPIR-V to LLVM.
void populateSPIRVToLLVMConversionPatterns(
- LLVMTypeConverter &typeConverter, RewritePatternSet &patterns,
+ const LLVMTypeConverter &typeConverter, RewritePatternSet &patterns,
spirv::ClientAPI clientAPIForAddressSpaceMapping =
spirv::ClientAPI::Unknown);
/// Populates the given list with patterns for function conversion from SPIR-V
/// to LLVM.
void populateSPIRVToLLVMFunctionConversionPatterns(
- LLVMTypeConverter &typeConverter, RewritePatternSet &patterns);
+ const LLVMTypeConverter &typeConverter, RewritePatternSet &patterns);
/// Populates the given patterns for module conversion from SPIR-V to LLVM.
void populateSPIRVToLLVMModuleConversionPatterns(
- LLVMTypeConverter &typeConverter, RewritePatternSet &patterns);
+ const LLVMTypeConverter &typeConverter, RewritePatternSet &patterns);
} // namespace mlir
diff --git a/mlir/include/mlir/Conversion/TensorToSPIRV/TensorToSPIRV.h b/mlir/include/mlir/Conversion/TensorToSPIRV/TensorToSPIRV.h
index fc122df23bc5ce..49cb862f780514 100644
--- a/mlir/include/mlir/Conversion/TensorToSPIRV/TensorToSPIRV.h
+++ b/mlir/include/mlir/Conversion/TensorToSPIRV/TensorToSPIRV.h
@@ -30,7 +30,7 @@ class SPIRVTypeConverter;
/// variables. SPIR-V consumers in GPU drivers may or may not optimize that
/// away. So this has implications over register pressure. Therefore, a
/// threshold is used to control when the patterns should kick in.
-void populateTensorToSPIRVPatterns(SPIRVTypeConverter &typeConverter,
+void populateTensorToSPIRVPatterns(const SPIRVTypeConverter &typeConverter,
int64_t byteCountThreshold,
RewritePatternSet &patterns);
diff --git a/mlir/include/mlir/Conversion/TosaToLinalg/TosaToLinalg.h b/mlir/include/mlir/Conversion/TosaToLinalg/TosaToLinalg.h
index c84e4f17c38d88..192583f347b8a4 100644
--- a/mlir/include/mlir/Conversion/TosaToLinalg/TosaToLinalg.h
+++ b/mlir/include/mlir/Conversion/TosaToLinalg/TosaToLinalg.h
@@ -47,7 +47,7 @@ void addTosaToLinalgPasses(
void registerTosaToLinalgPipelines();
/// Populates conversion passes from TOSA dialect to Linalg dialect.
-void populateTosaToLinalgConversionPatterns(TypeConverter &converter,
+void populateTosaToLinalgConversionPatterns(const TypeConverter &converter,
RewritePatternSet *patterns);
/// Populates conversion passes from TOSA dialect to Linalg named operations.
diff --git a/mlir/include/mlir/Conversion/TosaToTensor/TosaToTensor.h b/mlir/include/mlir/Conversion/TosaToTensor/TosaToTensor.h
index 76a4b1b1563366..66c381f5174183 100644
--- a/mlir/include/mlir/Conversion/TosaToTensor/TosaToTensor.h
+++ b/mlir/include/mlir/Conversion/TosaToTensor/TosaToTensor.h
@@ -25,7 +25,7 @@ namespace tosa {
std::unique_ptr<Pass> createTosaToTensor();
-void populateTosaToTensorConversionPatterns(TypeConverter &converter,
+void populateTosaToTensorConversionPatterns(const TypeConverter &converter,
RewritePatternSet *patterns);
} // namespace tosa
diff --git a/mlir/include/mlir/Conversion/UBToLLVM/UBToLLVM.h b/mlir/include/mlir/Conversion/UBToLLVM/UBToLLVM.h
index 787f69bdc31e95..ff9cd4759b9e89 100644
--- a/mlir/include/mlir/Conversion/UBToLLVM/UBToLLVM.h
+++ b/mlir/include/mlir/Conversion/UBToLLVM/UBToLLVM.h
@@ -22,7 +22,7 @@ class Pass;
#include "mlir/Conversion/Passes.h.inc"
namespace ub {
-void populateUBToLLVMConversionPatterns(LLVMTypeConverter &converter,
+void populateUBToLLVMConversionPatterns(const LLVMTypeConverter &converter,
RewritePatternSet &patterns);
void registerConvertUBToLLVMInterface(DialectRegistry ®istry);
diff --git a/mlir/include/mlir/Conversion/UBToSPIRV/UBToSPIRV.h b/mlir/include/mlir/Conversion/UBToSPIRV/UBToSPIRV.h
index 3843f2707a520d..5e6cf7fac6a220 100644
--- a/mlir/include/mlir/Conversion/UBToSPIRV/UBToSPIRV.h
+++ b/mlir/include/mlir/Conversion/UBToSPIRV/UBToSPIRV.h
@@ -21,7 +21,7 @@ class Pass;
#include "mlir/Conversion/Passes.h.inc"
namespace ub {
-void populateUBToSPIRVConversionPatterns(SPIRVTypeConverter &converter,
+void populateUBToSPIRVConversionPatterns(const SPIRVTypeConverter &converter,
RewritePatternSet &patterns);
} // namespace ub
} // namespace mlir
diff --git a/mlir/include/mlir/Conversion/VectorToLLVM/ConvertVectorToLLVM.h b/mlir/include/mlir/Conversion/VectorToLLVM/ConvertVectorToLLVM.h
index ecd33779236cc3..5fda62e3584c79 100644
--- a/mlir/include/mlir/Conversion/VectorToLLVM/ConvertVectorToLLVM.h
+++ b/mlir/include/mlir/Conversion/VectorToLLVM/ConvertVectorToLLVM.h
@@ -16,12 +16,12 @@ class LLVMTypeConverter;
/// Collect a set of patterns to convert from Vector contractions to LLVM Matrix
/// Intrinsics. To lower to assembly, the LLVM flag -lower-matrix-intrinsics
/// will be needed when invoking LLVM.
-void populateVectorToLLVMMatrixConversionPatterns(LLVMTypeConverter &converter,
- RewritePatternSet &patterns);
+void populateVectorToLLVMMatrixConversionPatterns(
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns);
/// Collect a set of patterns to convert from the Vector dialect to LLVM.
void populateVectorToLLVMConversionPatterns(
- LLVMTypeConverter &converter, RewritePatternSet &patterns,
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns,
bool reassociateFPReductions = false, bool force32BitVectorIndices = false);
} // namespace mlir
diff --git a/mlir/include/mlir/Conversion/VectorToSPIRV/VectorToSPIRV.h b/mlir/include/mlir/Conversion/VectorToSPIRV/VectorToSPIRV.h
index f8c02c54066b81..b64be4f733ec07 100644
--- a/mlir/include/mlir/Conversion/VectorToSPIRV/VectorToSPIRV.h
+++ b/mlir/include/mlir/Conversion/VectorToSPIRV/VectorToSPIRV.h
@@ -20,7 +20,7 @@ class SPIRVTypeConverter;
/// Appends to a pattern list additional patterns for translating Vector Ops to
/// SPIR-V ops.
-void populateVectorToSPIRVPatterns(SPIRVTypeConverter &typeConverter,
+void populateVectorToSPIRVPatterns(const SPIRVTypeConverter &typeConverter,
RewritePatternSet &patterns);
/// Appends patterns to convert vector reduction of the form:
diff --git a/mlir/include/mlir/Dialect/AMX/Transforms.h b/mlir/include/mlir/Dialect/AMX/Transforms.h
index 19608dff6f160f..d00ac52e274f9f 100644
--- a/mlir/include/mlir/Dialect/AMX/Transforms.h
+++ b/mlir/include/mlir/Dialect/AMX/Transforms.h
@@ -17,8 +17,8 @@ class RewritePatternSet;
/// Collect a set of patterns to lower AMX ops to ops that map to LLVM
/// intrinsics.
-void populateAMXLegalizeForLLVMExportPatterns(LLVMTypeConverter &converter,
- RewritePatternSet &patterns);
+void populateAMXLegalizeForLLVMExportPatterns(
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns);
/// Configure the target to support lowering AMX ops to ops that map to LLVM
/// intrinsics.
diff --git a/mlir/include/mlir/Dialect/Arith/Transforms/Passes.h b/mlir/include/mlir/Dialect/Arith/Transforms/Passes.h
index 9dc262cc72ed00..aee64475171a43 100644
--- a/mlir/include/mlir/Dialect/Arith/Transforms/Passes.h
+++ b/mlir/include/mlir/Dialect/Arith/Transforms/Passes.h
@@ -28,13 +28,15 @@ class NarrowTypeEmulationConverter;
/// types into supported ones. This is done by splitting original power-of-two
/// i2N integer types into two iN halves.
void populateArithWideIntEmulationPatterns(
- WideIntEmulationConverter &typeConverter, RewritePatternSet &patterns);
+ const WideIntEmulationConverter &typeConverter,
+ RewritePatternSet &patterns);
/// Adds patterns to emulate narrow Arith and Function ops into wide
/// supported types. Users need to add conversions about the computation
/// domain of narrow types.
void populateArithNarrowTypeEmulationPatterns(
- NarrowTypeEmulationConverter &typeConverter, RewritePatternSet &patterns);
+ const NarrowTypeEmulationConverter &typeConverter,
+ RewritePatternSet &patterns);
/// Populate the type conversions needed to emulate the unsupported
/// `sourceTypes` with `destType`
@@ -45,12 +47,12 @@ void populateEmulateUnsupportedFloatsConversions(TypeConverter &converter,
/// Add rewrite patterns for converting operations that use illegal float types
/// to ones that use legal ones.
void populateEmulateUnsupportedFloatsPatterns(RewritePatternSet &patterns,
- TypeConverter &converter);
+ const TypeConverter &converter);
/// Set up a dialect conversion to reject arithmetic operations on unsupported
/// float types.
void populateEmulateUnsupportedFloatsLegality(ConversionTarget &target,
- TypeConverter &converter);
+ const TypeConverter &converter);
/// Add patterns to expand Arith ceil/floor division ops.
void populateCeilFloorDivExpandOpsPatterns(RewritePatternSet &patterns);
diff --git a/mlir/include/mlir/Dialect/ArmSVE/Transforms/Transforms.h b/mlir/include/mlir/Dialect/ArmSVE/Transforms/Transforms.h
index dcc485789875d3..8665c8224cc45a 100644
--- a/mlir/include/mlir/Dialect/ArmSVE/Transforms/Transforms.h
+++ b/mlir/include/mlir/Dialect/ArmSVE/Transforms/Transforms.h
@@ -17,8 +17,8 @@ class RewritePatternSet;
/// Collect a set of patterns to lower ArmSVE ops to ops that map to LLVM
/// intrinsics.
-void populateArmSVELegalizeForLLVMExportPatterns(LLVMTypeConverter &converter,
- RewritePatternSet &patterns);
+void populateArmSVELegalizeForLLVMExportPatterns(
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns);
/// Configure the target to support lowering ArmSVE ops to ops that map to LLVM
/// intrinsics.
diff --git a/mlir/include/mlir/Dialect/Bufferization/Transforms/Bufferize.h b/mlir/include/mlir/Dialect/Bufferization/Transforms/Bufferize.h
index cab997e1aff297..1603dfcbae5589 100644
--- a/mlir/include/mlir/Dialect/Bufferization/Transforms/Bufferize.h
+++ b/mlir/include/mlir/Dialect/Bufferization/Transforms/Bufferize.h
@@ -60,7 +60,7 @@ void populateBufferizeMaterializationLegality(ConversionTarget &target);
///
/// In particular, these are the tensor_load/buffer_cast ops.
void populateEliminateBufferizeMaterializationsPatterns(
- BufferizeTypeConverter &typeConverter, RewritePatternSet &patterns);
+ const BufferizeTypeConverter &typeConverter, RewritePatternSet &patterns);
/// Bufferize `op` and its nested ops that implement `BufferizableOpInterface`.
///
diff --git a/mlir/include/mlir/Dialect/Func/Transforms/DecomposeCallGraphTypes.h b/mlir/include/mlir/Dialect/Func/Transforms/DecomposeCallGraphTypes.h
index 4af4c65148da3d..1d311b37b37a4f 100644
--- a/mlir/include/mlir/Dialect/Func/Transforms/DecomposeCallGraphTypes.h
+++ b/mlir/include/mlir/Dialect/Func/Transforms/DecomposeCallGraphTypes.h
@@ -85,7 +85,7 @@ class ValueDecomposer {
/// Populates the patterns needed to drive the conversion process for
/// decomposing call graph types with the given `ValueDecomposer`.
void populateDecomposeCallGraphTypesPatterns(MLIRContext *context,
- TypeConverter &typeConverter,
+ const TypeConverter &typeConverter,
ValueDecomposer &decomposer,
RewritePatternSet &patterns);
diff --git a/mlir/include/mlir/Dialect/Func/Transforms/FuncConversions.h b/mlir/include/mlir/Dialect/Func/Transforms/FuncConversions.h
index adcd3087b7e2fd..642f704c144301 100644
--- a/mlir/include/mlir/Dialect/Func/Transforms/FuncConversions.h
+++ b/mlir/include/mlir/Dialect/Func/Transforms/FuncConversions.h
@@ -29,7 +29,7 @@ class RewritePatternSet;
/// Add a pattern to the given pattern list to convert the operand and result
/// types of a CallOp with the given type converter.
void populateCallOpTypeConversionPattern(RewritePatternSet &patterns,
- TypeConverter &converter);
+ const TypeConverter &converter);
/// Add a pattern to the given pattern list to rewrite branch operations to use
/// operands that have been legalized by the conversion framework. This can only
@@ -41,25 +41,25 @@ void populateCallOpTypeConversionPattern(RewritePatternSet &patterns,
/// shouldConvertBranchOperand. This callback should return true if branchOp's
/// operand at index idx should be converted.
void populateBranchOpInterfaceTypeConversionPattern(
- RewritePatternSet &patterns, TypeConverter &converter,
+ RewritePatternSet &patterns, const TypeConverter &converter,
function_ref<bool(BranchOpInterface branchOp, int idx)>
shouldConvertBranchOperand = nullptr);
/// Return true if op is a BranchOpInterface op whose operands are all legal
/// according to converter.
-bool isLegalForBranchOpInterfaceTypeConversionPattern(Operation *op,
- TypeConverter &converter);
+bool isLegalForBranchOpInterfaceTypeConversionPattern(
+ Operation *op, const TypeConverter &converter);
/// Add a pattern to the given pattern list to rewrite `return` ops to use
/// operands that have been legalized by the conversion framework.
void populateReturnOpTypeConversionPattern(RewritePatternSet &patterns,
- TypeConverter &converter);
+ const TypeConverter &converter);
/// For ReturnLike ops (except `return`), return True. If op is a `return` &&
/// returnOpAlwaysLegal is false, legalize op according to converter. Otherwise,
/// return false.
bool isLegalForReturnOpTypeConversionPattern(Operation *op,
- TypeConverter &converter,
+ const TypeConverter &converter,
bool returnOpAlwaysLegal = false);
/// Return true if op is neither BranchOpInterface nor ReturnLike.
diff --git a/mlir/include/mlir/Dialect/Func/Transforms/OneToNFuncConversions.h b/mlir/include/mlir/Dialect/Func/Transforms/OneToNFuncConversions.h
index 2fba342ea80e71..c9e407daf9bf8c 100644
--- a/mlir/include/mlir/Dialect/Func/Transforms/OneToNFuncConversions.h
+++ b/mlir/include/mlir/Dialect/Func/Transforms/OneToNFuncConversions.h
@@ -18,7 +18,7 @@ namespace mlir {
// Populates the provided pattern set with patterns that do 1:N type conversions
// on func ops. This is intended to be used with `applyPartialOneToNConversion`.
-void populateFuncTypeConversionPatterns(TypeConverter &typeConverter,
+void populateFuncTypeConversionPatterns(const TypeConverter &typeConverter,
RewritePatternSet &patterns);
} // namespace mlir
diff --git a/mlir/include/mlir/Dialect/Math/Transforms/Passes.h b/mlir/include/mlir/Dialect/Math/Transforms/Passes.h
index 2974bb344ad965..74ba91322ea857 100644
--- a/mlir/include/mlir/Dialect/Math/Transforms/Passes.h
+++ b/mlir/include/mlir/Dialect/Math/Transforms/Passes.h
@@ -62,7 +62,7 @@ void populateExtendToSupportedTypesTypeConverter(
void populateExtendToSupportedTypesConversionTarget(
ConversionTarget &target, TypeConverter &typeConverter);
void populateExtendToSupportedTypesPatterns(RewritePatternSet &patterns,
- TypeConverter &typeConverter);
+ const TypeConverter &typeConverter);
} // namespace math
} // namespace mlir
diff --git a/mlir/include/mlir/Dialect/MemRef/Transforms/Transforms.h b/mlir/include/mlir/Dialect/MemRef/Transforms/Transforms.h
index 4444a7651b4072..62a2297c80e788 100644
--- a/mlir/include/mlir/Dialect/MemRef/Transforms/Transforms.h
+++ b/mlir/include/mlir/Dialect/MemRef/Transforms/Transforms.h
@@ -72,7 +72,7 @@ void populateExpandReallocPatterns(RewritePatternSet &patterns,
/// Appends patterns for emulating wide integer memref operations with ops over
/// narrower integer types.
void populateMemRefWideIntEmulationPatterns(
- arith::WideIntEmulationConverter &typeConverter,
+ const arith::WideIntEmulationConverter &typeConverter,
RewritePatternSet &patterns);
/// Appends type conversions for emulating wide integer memref operations with
@@ -83,7 +83,7 @@ void populateMemRefWideIntEmulationConversions(
/// Appends patterns for emulating memref operations over narrow types with ops
/// over wider types.
void populateMemRefNarrowTypeEmulationPatterns(
- arith::NarrowTypeEmulationConverter &typeConverter,
+ const arith::NarrowTypeEmulationConverter &typeConverter,
RewritePatternSet &patterns);
/// Appends type conversions for emulating memref operations over narrow types
diff --git a/mlir/include/mlir/Dialect/SCF/Transforms/Patterns.h b/mlir/include/mlir/Dialect/SCF/Transforms/Patterns.h
index 5e66774d2f1430..b87407d302a823 100644
--- a/mlir/include/mlir/Dialect/SCF/Transforms/Patterns.h
+++ b/mlir/include/mlir/Dialect/SCF/Transforms/Patterns.h
@@ -50,12 +50,12 @@ class ForLoopPipeliningPattern : public OpRewritePattern<ForOp> {
/// corresponding scf.yield ops need to update their types accordingly to the
/// TypeConverter, but otherwise don't care what type conversions are happening.
void populateSCFStructuralTypeConversionsAndLegality(
- TypeConverter &typeConverter, RewritePatternSet &patterns,
+ const TypeConverter &typeConverter, RewritePatternSet &patterns,
ConversionTarget &target);
/// Similar to `populateSCFStructuralTypeConversionsAndLegality` but does not
/// populate the conversion target.
-void populateSCFStructuralTypeConversions(TypeConverter &typeConverter,
+void populateSCFStructuralTypeConversions(const TypeConverter &typeConverter,
RewritePatternSet &patterns);
/// Updates the ConversionTarget with dynamic legality of SCF operations based
@@ -66,8 +66,8 @@ void populateSCFStructuralTypeConversionTarget(
/// Populates the provided pattern set with patterns that do 1:N type
/// conversions on (some) SCF ops. This is intended to be used with
/// applyPartialOneToNConversion.
-void populateSCFStructuralOneToNTypeConversions(TypeConverter &typeConverter,
- RewritePatternSet &patterns);
+void populateSCFStructuralOneToNTypeConversions(
+ const TypeConverter &typeConverter, RewritePatternSet &patterns);
/// Populate patterns for SCF software pipelining transformation. See the
/// ForLoopPipeliningPattern for the transformation details.
diff --git a/mlir/include/mlir/Dialect/SPIRV/Transforms/SPIRVConversion.h b/mlir/include/mlir/Dialect/SPIRV/Transforms/SPIRVConversion.h
index f54c93a09e7270..bed4d66ccd6cbe 100644
--- a/mlir/include/mlir/Dialect/SPIRV/Transforms/SPIRVConversion.h
+++ b/mlir/include/mlir/Dialect/SPIRV/Transforms/SPIRVConversion.h
@@ -135,7 +135,7 @@ class SPIRVConversionTarget : public ConversionTarget {
/// `func` op to the SPIR-V dialect. These patterns do not handle shader
/// interface/ABI; they convert function parameters to be of SPIR-V allowed
/// types.
-void populateBuiltinFuncToSPIRVPatterns(SPIRVTypeConverter &typeConverter,
+void populateBuiltinFuncToSPIRVPatterns(const SPIRVTypeConverter &typeConverter,
RewritePatternSet &patterns);
void populateFuncOpVectorRewritePatterns(RewritePatternSet &patterns);
diff --git a/mlir/include/mlir/Dialect/SparseTensor/Transforms/Passes.h b/mlir/include/mlir/Dialect/SparseTensor/Transforms/Passes.h
index d22df6a7857c1d..6ccbc40bdd6034 100644
--- a/mlir/include/mlir/Dialect/SparseTensor/Transforms/Passes.h
+++ b/mlir/include/mlir/Dialect/SparseTensor/Transforms/Passes.h
@@ -154,7 +154,7 @@ struct SparseIterationTypeConverter : public OneToNTypeConverter {
SparseIterationTypeConverter();
};
-void populateLowerSparseIterationToSCFPatterns(TypeConverter &converter,
+void populateLowerSparseIterationToSCFPatterns(const TypeConverter &converter,
RewritePatternSet &patterns);
std::unique_ptr<Pass> createLowerSparseIterationToSCFPass();
@@ -170,7 +170,7 @@ class SparseTensorTypeToPtrConverter : public TypeConverter {
};
/// Sets up sparse tensor conversion rules.
-void populateSparseTensorConversionPatterns(TypeConverter &typeConverter,
+void populateSparseTensorConversionPatterns(const TypeConverter &typeConverter,
RewritePatternSet &patterns);
std::unique_ptr<Pass> createSparseTensorConversionPass();
@@ -186,7 +186,7 @@ class SparseTensorTypeToBufferConverter : public TypeConverter {
};
/// Sets up sparse tensor codegen rules.
-void populateSparseTensorCodegenPatterns(TypeConverter &typeConverter,
+void populateSparseTensorCodegenPatterns(const TypeConverter &typeConverter,
RewritePatternSet &patterns,
bool createSparseDeallocs,
bool enableBufferInitialization);
@@ -244,7 +244,7 @@ class StorageSpecifierToLLVMTypeConverter : public TypeConverter {
StorageSpecifierToLLVMTypeConverter();
};
-void populateStorageSpecifierToLLVMPatterns(TypeConverter &converter,
+void populateStorageSpecifierToLLVMPatterns(const TypeConverter &converter,
RewritePatternSet &patterns);
std::unique_ptr<Pass> createStorageSpecifierToLLVMPass();
diff --git a/mlir/include/mlir/Dialect/Vector/Transforms/VectorRewritePatterns.h b/mlir/include/mlir/Dialect/Vector/Transforms/VectorRewritePatterns.h
index 456d021bc79358..a59f06f3c1ef1b 100644
--- a/mlir/include/mlir/Dialect/Vector/Transforms/VectorRewritePatterns.h
+++ b/mlir/include/mlir/Dialect/Vector/Transforms/VectorRewritePatterns.h
@@ -366,7 +366,7 @@ void populateVectorMaskMaterializationPatterns(RewritePatternSet &patterns,
/// Appends patterns for emulating vector operations over narrow types with ops
/// over wider types.
void populateVectorNarrowTypeEmulationPatterns(
- arith::NarrowTypeEmulationConverter &typeConverter,
+ const arith::NarrowTypeEmulationConverter &typeConverter,
RewritePatternSet &patterns);
/// Rewrite a vector `bitcast(trunci)` to use a more efficient sequence of
@@ -403,10 +403,9 @@ void populateVectorLinearizeTypeConversionsAndLegality(
/// Populates patterns for linearizing ND (N >= 2) vector operations to 1D
/// vector shuffle operations.
-void populateVectorLinearizeShuffleLikeOpsPatterns(TypeConverter &typeConverter,
- RewritePatternSet &patterns,
- ConversionTarget &target,
- unsigned targetBitWidth);
+void populateVectorLinearizeShuffleLikeOpsPatterns(
+ const TypeConverter &typeConverter, RewritePatternSet &patterns,
+ ConversionTarget &target, unsigned targetBitWidth);
} // namespace vector
} // namespace mlir
diff --git a/mlir/include/mlir/Dialect/X86Vector/Transforms.h b/mlir/include/mlir/Dialect/X86Vector/Transforms.h
index fe98ebd955aa31..d54111ca41e69b 100644
--- a/mlir/include/mlir/Dialect/X86Vector/Transforms.h
+++ b/mlir/include/mlir/Dialect/X86Vector/Transforms.h
@@ -176,7 +176,7 @@ void populateSpecializedTransposeLoweringPatterns(
/// Collect a set of patterns to lower X86Vector ops to ops that map to LLVM
/// intrinsics.
void populateX86VectorLegalizeForLLVMExportPatterns(
- LLVMTypeConverter &converter, RewritePatternSet &patterns);
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns);
/// Configure the target to support lowering X86Vector ops to ops that map to
/// LLVM intrinsics.
diff --git a/mlir/include/mlir/Transforms/OneToNTypeConversion.h b/mlir/include/mlir/Transforms/OneToNTypeConversion.h
index 4c689ba219e884..c59a3a52f028f3 100644
--- a/mlir/include/mlir/Transforms/OneToNTypeConversion.h
+++ b/mlir/include/mlir/Transforms/OneToNTypeConversion.h
@@ -53,7 +53,7 @@ class OneToNTypeConverter : public TypeConverter {
/// `TypeConverter::convertSignatureArgs` and exists here with a different
/// name to reflect the broader semantic.
LogicalResult computeTypeMapping(TypeRange types,
- SignatureConversion &result) {
+ SignatureConversion &result) const {
return convertSignatureArgs(types, result);
}
@@ -126,24 +126,25 @@ class RewritePatternWithConverter : public mlir::RewritePattern {
/// Construct a conversion pattern with the given converter, and forward the
/// remaining arguments to RewritePattern.
template <typename... Args>
- RewritePatternWithConverter(TypeConverter &typeConverter, Args &&...args)
+ RewritePatternWithConverter(const TypeConverter &typeConverter,
+ Args &&...args)
: RewritePattern(std::forward<Args>(args)...),
typeConverter(&typeConverter) {}
/// Return the type converter held by this pattern, or nullptr if the pattern
/// does not require type conversion.
- TypeConverter *getTypeConverter() const { return typeConverter; }
+ const TypeConverter *getTypeConverter() const { return typeConverter; }
template <typename ConverterTy>
std::enable_if_t<std::is_base_of<TypeConverter, ConverterTy>::value,
- ConverterTy *>
+ const ConverterTy *>
getTypeConverter() const {
- return static_cast<ConverterTy *>(typeConverter);
+ return static_cast<const ConverterTy *>(typeConverter);
}
protected:
/// A type converter for use by this pattern.
- TypeConverter *const typeConverter;
+ const TypeConverter *const typeConverter;
};
/// Specialization of `PatternRewriter` that `OneToNConversionPattern`s use. The
@@ -212,8 +213,8 @@ class OneToNConversionPattern : public RewritePatternWithConverter {
template <typename SourceOp>
class OneToNOpConversionPattern : public OneToNConversionPattern {
public:
- OneToNOpConversionPattern(TypeConverter &typeConverter, MLIRContext *context,
- PatternBenefit benefit = 1,
+ OneToNOpConversionPattern(const TypeConverter &typeConverter,
+ MLIRContext *context, PatternBenefit benefit = 1,
ArrayRef<StringRef> generatedNames = {})
: OneToNConversionPattern(typeConverter, SourceOp::getOperationName(),
benefit, context, generatedNames) {}
@@ -302,11 +303,11 @@ applyPartialOneToNConversion(Operation *op, OneToNTypeConverter &typeConverter,
/// ops which use FunctionType to represent their type. This is intended to be
/// used with the 1:N dialect conversion.
void populateOneToNFunctionOpInterfaceTypeConversionPattern(
- StringRef functionLikeOpName, TypeConverter &converter,
+ StringRef functionLikeOpName, const TypeConverter &converter,
RewritePatternSet &patterns);
template <typename FuncOpT>
void populateOneToNFunctionOpInterfaceTypeConversionPattern(
- TypeConverter &converter, RewritePatternSet &patterns) {
+ const TypeConverter &converter, RewritePatternSet &patterns) {
populateOneToNFunctionOpInterfaceTypeConversionPattern(
FuncOpT::getOperationName(), converter, patterns);
}
diff --git a/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp b/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp
index 7112d1607dfdca..2b33f3773dc7d1 100644
--- a/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp
+++ b/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp
@@ -277,7 +277,7 @@ struct RawBufferOpLowering : public ConvertOpToLLVMPattern<GpuOp> {
};
struct LDSBarrierOpLowering : public ConvertOpToLLVMPattern<LDSBarrierOp> {
- LDSBarrierOpLowering(LLVMTypeConverter &converter, Chipset chipset)
+ LDSBarrierOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
: ConvertOpToLLVMPattern<LDSBarrierOp>(converter), chipset(chipset) {}
Chipset chipset;
@@ -335,7 +335,7 @@ struct LDSBarrierOpLowering : public ConvertOpToLLVMPattern<LDSBarrierOp> {
};
struct SchedBarrierOpLowering : public ConvertOpToLLVMPattern<SchedBarrierOp> {
- SchedBarrierOpLowering(LLVMTypeConverter &converter, Chipset chipset)
+ SchedBarrierOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
: ConvertOpToLLVMPattern<SchedBarrierOp>(converter), chipset(chipset) {}
Chipset chipset;
@@ -725,7 +725,7 @@ struct WMMAOpLowering : public ConvertOpToLLVMPattern<WMMAOp> {
namespace {
struct ExtPackedFp8OpLowering final
: public ConvertOpToLLVMPattern<ExtPackedFp8Op> {
- ExtPackedFp8OpLowering(LLVMTypeConverter &converter, Chipset chipset)
+ ExtPackedFp8OpLowering(const LLVMTypeConverter &converter, Chipset chipset)
: ConvertOpToLLVMPattern<amdgpu::ExtPackedFp8Op>(converter),
chipset(chipset) {}
Chipset chipset;
@@ -737,7 +737,8 @@ struct ExtPackedFp8OpLowering final
struct PackedTrunc2xFp8OpLowering final
: public ConvertOpToLLVMPattern<PackedTrunc2xFp8Op> {
- PackedTrunc2xFp8OpLowering(LLVMTypeConverter &converter, Chipset chipset)
+ PackedTrunc2xFp8OpLowering(const LLVMTypeConverter &converter,
+ Chipset chipset)
: ConvertOpToLLVMPattern<amdgpu::PackedTrunc2xFp8Op>(converter),
chipset(chipset) {}
Chipset chipset;
@@ -749,7 +750,8 @@ struct PackedTrunc2xFp8OpLowering final
struct PackedStochRoundFp8OpLowering final
: public ConvertOpToLLVMPattern<PackedStochRoundFp8Op> {
- PackedStochRoundFp8OpLowering(LLVMTypeConverter &converter, Chipset chipset)
+ PackedStochRoundFp8OpLowering(const LLVMTypeConverter &converter,
+ Chipset chipset)
: ConvertOpToLLVMPattern<amdgpu::PackedStochRoundFp8Op>(converter),
chipset(chipset) {}
Chipset chipset;
@@ -880,7 +882,7 @@ LogicalResult PackedStochRoundFp8OpLowering::matchAndRewrite(
// Implement the AMDGPU_DPPLowering class that will convert the amdgpu.dpp
// operation into the corresponding ROCDL instructions.
struct AMDGPUDPPLowering : public ConvertOpToLLVMPattern<DPPOp> {
- AMDGPUDPPLowering(LLVMTypeConverter &converter, Chipset chipset)
+ AMDGPUDPPLowering(const LLVMTypeConverter &converter, Chipset chipset)
: ConvertOpToLLVMPattern<DPPOp>(converter), chipset(chipset) {}
Chipset chipset;
@@ -1052,9 +1054,9 @@ struct ConvertAMDGPUToROCDLPass
};
} // namespace
-void mlir::populateAMDGPUToROCDLConversionPatterns(LLVMTypeConverter &converter,
- RewritePatternSet &patterns,
- Chipset chipset) {
+void mlir::populateAMDGPUToROCDLConversionPatterns(
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns,
+ Chipset chipset) {
patterns
.add<RawBufferOpLowering<RawBufferLoadOp, ROCDL::RawPtrBufferLoadOp>,
RawBufferOpLowering<RawBufferStoreOp, ROCDL::RawPtrBufferStoreOp>,
diff --git a/mlir/lib/Conversion/ArithToLLVM/ArithToLLVM.cpp b/mlir/lib/Conversion/ArithToLLVM/ArithToLLVM.cpp
index d882f1184f4570..aac24f113d891f 100644
--- a/mlir/lib/Conversion/ArithToLLVM/ArithToLLVM.cpp
+++ b/mlir/lib/Conversion/ArithToLLVM/ArithToLLVM.cpp
@@ -520,7 +520,7 @@ void mlir::arith::registerConvertArithToLLVMInterface(
//===----------------------------------------------------------------------===//
void mlir::arith::populateArithToLLVMConversionPatterns(
- LLVMTypeConverter &converter, RewritePatternSet &patterns) {
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns) {
// clang-format off
patterns.add<
AddFOpLowering,
diff --git a/mlir/lib/Conversion/ArithToSPIRV/ArithToSPIRV.cpp b/mlir/lib/Conversion/ArithToSPIRV/ArithToSPIRV.cpp
index e6c01f063e8b80..382b93cd9f21b1 100644
--- a/mlir/lib/Conversion/ArithToSPIRV/ArithToSPIRV.cpp
+++ b/mlir/lib/Conversion/ArithToSPIRV/ArithToSPIRV.cpp
@@ -1262,7 +1262,7 @@ class MinNumMaxNumFOpPattern final : public OpConversionPattern<Op> {
//===----------------------------------------------------------------------===//
void mlir::arith::populateArithToSPIRVPatterns(
- SPIRVTypeConverter &typeConverter, RewritePatternSet &patterns) {
+ const SPIRVTypeConverter &typeConverter, RewritePatternSet &patterns) {
// clang-format off
patterns.add<
ConstantCompositeOpPattern,
diff --git a/mlir/lib/Conversion/ComplexToLLVM/ComplexToLLVM.cpp b/mlir/lib/Conversion/ComplexToLLVM/ComplexToLLVM.cpp
index 0a3c3a330ff699..fcc18e6eee9980 100644
--- a/mlir/lib/Conversion/ComplexToLLVM/ComplexToLLVM.cpp
+++ b/mlir/lib/Conversion/ComplexToLLVM/ComplexToLLVM.cpp
@@ -323,7 +323,7 @@ struct SubOpConversion : public ConvertOpToLLVMPattern<complex::SubOp> {
} // namespace
void mlir::populateComplexToLLVMConversionPatterns(
- LLVMTypeConverter &converter, RewritePatternSet &patterns) {
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns) {
// clang-format off
patterns.add<
AbsOpConversion,
diff --git a/mlir/lib/Conversion/ComplexToSPIRV/ComplexToSPIRV.cpp b/mlir/lib/Conversion/ComplexToSPIRV/ComplexToSPIRV.cpp
index d531659e0623a0..5e7d2d8491533f 100644
--- a/mlir/lib/Conversion/ComplexToSPIRV/ComplexToSPIRV.cpp
+++ b/mlir/lib/Conversion/ComplexToSPIRV/ComplexToSPIRV.cpp
@@ -102,8 +102,8 @@ struct ImOpPattern final : OpConversionPattern<complex::ImOp> {
// Pattern population
//===----------------------------------------------------------------------===//
-void mlir::populateComplexToSPIRVPatterns(SPIRVTypeConverter &typeConverter,
- RewritePatternSet &patterns) {
+void mlir::populateComplexToSPIRVPatterns(
+ const SPIRVTypeConverter &typeConverter, RewritePatternSet &patterns) {
MLIRContext *context = patterns.getContext();
patterns.add<ConstantOpPattern, CreateOpPattern, ReOpPattern, ImOpPattern>(
diff --git a/mlir/lib/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.cpp b/mlir/lib/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.cpp
index b8e5aec25286d2..e5c735e10703a7 100644
--- a/mlir/lib/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.cpp
+++ b/mlir/lib/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.cpp
@@ -43,7 +43,7 @@ namespace {
/// ignored by the default lowering but should be propagated by any custom
/// lowering.
struct AssertOpLowering : public ConvertOpToLLVMPattern<cf::AssertOp> {
- explicit AssertOpLowering(LLVMTypeConverter &typeConverter,
+ explicit AssertOpLowering(const LLVMTypeConverter &typeConverter,
bool abortOnFailedAssert = true)
: ConvertOpToLLVMPattern<cf::AssertOp>(typeConverter, /*benefit=*/1),
abortOnFailedAssert(abortOnFailedAssert) {}
@@ -201,7 +201,7 @@ struct SwitchOpLowering : public ConvertOpToLLVMPattern<cf::SwitchOp> {
} // namespace
void mlir::cf::populateControlFlowToLLVMConversionPatterns(
- LLVMTypeConverter &converter, RewritePatternSet &patterns) {
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns) {
// clang-format off
patterns.add<
AssertOpLowering,
@@ -212,7 +212,7 @@ void mlir::cf::populateControlFlowToLLVMConversionPatterns(
}
void mlir::cf::populateAssertToLLVMConversionPattern(
- LLVMTypeConverter &converter, RewritePatternSet &patterns,
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns,
bool abortOnFailure) {
patterns.add<AssertOpLowering>(converter, abortOnFailure);
}
diff --git a/mlir/lib/Conversion/ControlFlowToSPIRV/ControlFlowToSPIRV.cpp b/mlir/lib/Conversion/ControlFlowToSPIRV/ControlFlowToSPIRV.cpp
index f96bfd6f788b9c..d8f3e995109538 100644
--- a/mlir/lib/Conversion/ControlFlowToSPIRV/ControlFlowToSPIRV.cpp
+++ b/mlir/lib/Conversion/ControlFlowToSPIRV/ControlFlowToSPIRV.cpp
@@ -109,7 +109,7 @@ struct CondBranchOpPattern final : OpConversionPattern<cf::CondBranchOp> {
//===----------------------------------------------------------------------===//
void mlir::cf::populateControlFlowToSPIRVPatterns(
- SPIRVTypeConverter &typeConverter, RewritePatternSet &patterns) {
+ const SPIRVTypeConverter &typeConverter, RewritePatternSet &patterns) {
MLIRContext *context = patterns.getContext();
patterns.add<BranchOpPattern, CondBranchOpPattern>(typeConverter, context);
diff --git a/mlir/lib/Conversion/ConvertToSPIRV/ConvertToSPIRVPass.cpp b/mlir/lib/Conversion/ConvertToSPIRV/ConvertToSPIRVPass.cpp
index f85993e52b4052..4b7f7ff114deeb 100644
--- a/mlir/lib/Conversion/ConvertToSPIRV/ConvertToSPIRVPass.cpp
+++ b/mlir/lib/Conversion/ConvertToSPIRV/ConvertToSPIRVPass.cpp
@@ -55,7 +55,7 @@ void mapToMemRef(Operation *op, spirv::TargetEnvAttr &targetAttr) {
}
/// Populate patterns for each dialect.
-void populateConvertToSPIRVPatterns(SPIRVTypeConverter &typeConverter,
+void populateConvertToSPIRVPatterns(const SPIRVTypeConverter &typeConverter,
ScfToSPIRVContext &scfToSPIRVContext,
RewritePatternSet &patterns) {
arith::populateCeilFloorDivExpandOpsPatterns(patterns);
diff --git a/mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp b/mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp
index 2cc77e8fd41b9a..27c43e0daad072 100644
--- a/mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp
+++ b/mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp
@@ -722,12 +722,12 @@ struct ReturnOpLowering : public ConvertOpToLLVMPattern<func::ReturnOp> {
} // namespace
void mlir::populateFuncToLLVMFuncOpConversionPattern(
- LLVMTypeConverter &converter, RewritePatternSet &patterns) {
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns) {
patterns.add<FuncOpConversion>(converter);
}
void mlir::populateFuncToLLVMConversionPatterns(
- LLVMTypeConverter &converter, RewritePatternSet &patterns,
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns,
const SymbolTable *symbolTable) {
populateFuncToLLVMFuncOpConversionPattern(converter, patterns);
patterns.add<CallIndirectOpLowering>(converter);
diff --git a/mlir/lib/Conversion/FuncToSPIRV/FuncToSPIRV.cpp b/mlir/lib/Conversion/FuncToSPIRV/FuncToSPIRV.cpp
index 4740b7cc6c385b..a5c2b92ca191df 100644
--- a/mlir/lib/Conversion/FuncToSPIRV/FuncToSPIRV.cpp
+++ b/mlir/lib/Conversion/FuncToSPIRV/FuncToSPIRV.cpp
@@ -87,7 +87,7 @@ class CallOpPattern final : public OpConversionPattern<func::CallOp> {
// Pattern population
//===----------------------------------------------------------------------===//
-void mlir::populateFuncToSPIRVPatterns(SPIRVTypeConverter &typeConverter,
+void mlir::populateFuncToSPIRVPatterns(const SPIRVTypeConverter &typeConverter,
RewritePatternSet &patterns) {
MLIRContext *context = patterns.getContext();
diff --git a/mlir/lib/Conversion/GPUCommon/IndexIntrinsicsOpLowering.h b/mlir/lib/Conversion/GPUCommon/IndexIntrinsicsOpLowering.h
index eaf1554a83f89a..ba4de3cb6b3e92 100644
--- a/mlir/lib/Conversion/GPUCommon/IndexIntrinsicsOpLowering.h
+++ b/mlir/lib/Conversion/GPUCommon/IndexIntrinsicsOpLowering.h
@@ -36,12 +36,12 @@ struct OpLowering : public ConvertOpToLLVMPattern<Op> {
IntrType intrType;
public:
- explicit OpLowering(LLVMTypeConverter &typeConverter)
+ explicit OpLowering(const LLVMTypeConverter &typeConverter)
: ConvertOpToLLVMPattern<Op>(typeConverter),
indexBitwidth(typeConverter.getIndexTypeBitwidth()),
indexKind(IndexKind::Other), intrType(IntrType::None) {}
- explicit OpLowering(LLVMTypeConverter &typeConverter, IndexKind indexKind,
+ explicit OpLowering(const LLVMTypeConverter &typeConverter, IndexKind indexKind,
IntrType intrType)
: ConvertOpToLLVMPattern<Op>(typeConverter),
indexBitwidth(typeConverter.getIndexTypeBitwidth()),
diff --git a/mlir/lib/Conversion/GPUCommon/OpToFuncCallLowering.h b/mlir/lib/Conversion/GPUCommon/OpToFuncCallLowering.h
index 8ff4d4ec67b9fd..edfbb7895ec378 100644
--- a/mlir/lib/Conversion/GPUCommon/OpToFuncCallLowering.h
+++ b/mlir/lib/Conversion/GPUCommon/OpToFuncCallLowering.h
@@ -42,7 +42,7 @@ namespace mlir {
template <typename SourceOp>
struct OpToFuncCallLowering : public ConvertOpToLLVMPattern<SourceOp> {
public:
- explicit OpToFuncCallLowering(LLVMTypeConverter &lowering, StringRef f32Func,
+ explicit OpToFuncCallLowering(const LLVMTypeConverter &lowering, StringRef f32Func,
StringRef f64Func, StringRef f32ApproxFunc,
StringRef f16Func)
: ConvertOpToLLVMPattern<SourceOp>(lowering), f32Func(f32Func),
diff --git a/mlir/lib/Conversion/GPUToLLVMSPV/GPUToLLVMSPV.cpp b/mlir/lib/Conversion/GPUToLLVMSPV/GPUToLLVMSPV.cpp
index 739a34e0aa610e..544f1f4a4f6a79 100644
--- a/mlir/lib/Conversion/GPUToLLVMSPV/GPUToLLVMSPV.cpp
+++ b/mlir/lib/Conversion/GPUToLLVMSPV/GPUToLLVMSPV.cpp
@@ -412,8 +412,8 @@ gpuAddressSpaceToOCLAddressSpace(gpu::AddressSpace addressSpace) {
}
} // namespace
-void populateGpuToLLVMSPVConversionPatterns(LLVMTypeConverter &typeConverter,
- RewritePatternSet &patterns) {
+void populateGpuToLLVMSPVConversionPatterns(
+ const LLVMTypeConverter &typeConverter, RewritePatternSet &patterns) {
patterns.add<GPUBarrierConversion, GPUReturnOpLowering, GPUShuffleConversion,
GPUSubgroupOpConversion<gpu::LaneIdOp>,
GPUSubgroupOpConversion<gpu::NumSubgroupsOp>,
diff --git a/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp b/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp
index 2b91a6c28c05e8..e83574b7342725 100644
--- a/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp
+++ b/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp
@@ -333,7 +333,7 @@ void mlir::configureGpuToNVVMConversionLegality(ConversionTarget &target) {
}
template <typename OpTy>
-static void populateOpPatterns(LLVMTypeConverter &converter,
+static void populateOpPatterns(const LLVMTypeConverter &converter,
RewritePatternSet &patterns, StringRef f32Func,
StringRef f64Func, StringRef f32ApproxFunc = "",
StringRef f16Func = "") {
@@ -343,12 +343,12 @@ static void populateOpPatterns(LLVMTypeConverter &converter,
}
void mlir::populateGpuSubgroupReduceOpLoweringPattern(
- LLVMTypeConverter &converter, RewritePatternSet &patterns) {
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns) {
patterns.add<GPUSubgroupReduceOpLowering>(converter);
}
-void mlir::populateGpuToNVVMConversionPatterns(LLVMTypeConverter &converter,
- RewritePatternSet &patterns) {
+void mlir::populateGpuToNVVMConversionPatterns(
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns) {
using gpu::index_lowering::IndexKind;
using gpu::index_lowering::IntrType;
populateWithGenerated(patterns);
diff --git a/mlir/lib/Conversion/GPUToNVVM/WmmaOpsToNvvm.cpp b/mlir/lib/Conversion/GPUToNVVM/WmmaOpsToNvvm.cpp
index b7fd454c60902f..2b040fddac7486 100644
--- a/mlir/lib/Conversion/GPUToNVVM/WmmaOpsToNvvm.cpp
+++ b/mlir/lib/Conversion/GPUToNVVM/WmmaOpsToNvvm.cpp
@@ -388,7 +388,7 @@ LLVM::LLVMStructType mlir::convertMMAToLLVMType(gpu::MMAMatrixType type) {
}
void mlir::populateGpuWMMAToNVVMConversionPatterns(
- LLVMTypeConverter &converter, RewritePatternSet &patterns) {
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns) {
patterns.add<WmmaLoadOpToNVVMLowering, WmmaMmaOpToNVVMLowering,
WmmaStoreOpToNVVMLowering, WmmaConstantOpToNVVMLowering,
WmmaElementwiseOpToNVVMLowering>(converter);
diff --git a/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp b/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp
index 482c9e2c2d0017..e253037e0edcea 100644
--- a/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp
+++ b/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp
@@ -343,7 +343,7 @@ void mlir::configureGpuToROCDLConversionLegality(ConversionTarget &target) {
}
template <typename OpTy>
-static void populateOpPatterns(LLVMTypeConverter &converter,
+static void populateOpPatterns(const LLVMTypeConverter &converter,
RewritePatternSet &patterns, StringRef f32Func,
StringRef f64Func, StringRef f32ApproxFunc,
StringRef f16Func) {
@@ -353,7 +353,7 @@ static void populateOpPatterns(LLVMTypeConverter &converter,
}
void mlir::populateGpuToROCDLConversionPatterns(
- LLVMTypeConverter &converter, RewritePatternSet &patterns,
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns,
mlir::gpu::amd::Runtime runtime) {
using gpu::index_lowering::IndexKind;
using gpu::index_lowering::IntrType;
diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
index 53b4c720ae56d2..3cc64b82950b56 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
@@ -59,7 +59,8 @@ class SingleDimLaunchConfigConversion : public OpConversionPattern<SourceOp> {
/// attribute on the surrounding FuncOp is used to replace the gpu::BlockDimOp.
class WorkGroupSizeConversion : public OpConversionPattern<gpu::BlockDimOp> {
public:
- WorkGroupSizeConversion(TypeConverter &typeConverter, MLIRContext *context)
+ WorkGroupSizeConversion(const TypeConverter &typeConverter,
+ MLIRContext *context)
: OpConversionPattern(typeConverter, context, /*benefit*/ 10) {}
LogicalResult
@@ -728,7 +729,7 @@ LogicalResult GPUPrintfConversion::matchAndRewrite(
// GPU To SPIRV Patterns.
//===----------------------------------------------------------------------===//
-void mlir::populateGPUToSPIRVPatterns(SPIRVTypeConverter &typeConverter,
+void mlir::populateGPUToSPIRVPatterns(const SPIRVTypeConverter &typeConverter,
RewritePatternSet &patterns) {
patterns.add<
GPUBarrierConversion, GPUFuncOpConversion, GPUModuleConversion,
diff --git a/mlir/lib/Conversion/GPUToSPIRV/WmmaOpsToSPIRV.cpp b/mlir/lib/Conversion/GPUToSPIRV/WmmaOpsToSPIRV.cpp
index 92cc0eadb9784c..1b0f0235278916 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/WmmaOpsToSPIRV.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/WmmaOpsToSPIRV.cpp
@@ -293,7 +293,7 @@ struct WmmaMmaOpToSPIRVLowering final
} // namespace mlir
void mlir::populateGpuWMMAToSPIRVCoopMatrixKHRConversionPatterns(
- SPIRVTypeConverter &converter, RewritePatternSet &patterns) {
+ const SPIRVTypeConverter &converter, RewritePatternSet &patterns) {
using namespace mlir;
MLIRContext *context = patterns.getContext();
patterns.add<khr::WmmaLoadOpToSPIRVLowering, khr::WmmaMmaOpToSPIRVLowering,
diff --git a/mlir/lib/Conversion/IndexToLLVM/IndexToLLVM.cpp b/mlir/lib/Conversion/IndexToLLVM/IndexToLLVM.cpp
index 9d8a5d8a0e1c06..0473bb59fa6aa3 100644
--- a/mlir/lib/Conversion/IndexToLLVM/IndexToLLVM.cpp
+++ b/mlir/lib/Conversion/IndexToLLVM/IndexToLLVM.cpp
@@ -292,7 +292,7 @@ using ConvertIndexBoolConstant =
//===----------------------------------------------------------------------===//
void index::populateIndexToLLVMConversionPatterns(
- LLVMTypeConverter &typeConverter, RewritePatternSet &patterns) {
+ const LLVMTypeConverter &typeConverter, RewritePatternSet &patterns) {
patterns.insert<
// clang-format off
ConvertIndexAdd,
diff --git a/mlir/lib/Conversion/IndexToSPIRV/IndexToSPIRV.cpp b/mlir/lib/Conversion/IndexToSPIRV/IndexToSPIRV.cpp
index b58efc096e2eaf..b4cc8324883eeb 100644
--- a/mlir/lib/Conversion/IndexToSPIRV/IndexToSPIRV.cpp
+++ b/mlir/lib/Conversion/IndexToSPIRV/IndexToSPIRV.cpp
@@ -338,8 +338,8 @@ struct ConvertIndexSizeOf final : OpConversionPattern<SizeOfOp> {
// Pattern Population
//===----------------------------------------------------------------------===//
-void index::populateIndexToSPIRVPatterns(SPIRVTypeConverter &typeConverter,
- RewritePatternSet &patterns) {
+void index::populateIndexToSPIRVPatterns(
+ const SPIRVTypeConverter &typeConverter, RewritePatternSet &patterns) {
patterns.add<
// clang-format off
ConvertIndexAdd,
diff --git a/mlir/lib/Conversion/MathToLLVM/MathToLLVM.cpp b/mlir/lib/Conversion/MathToLLVM/MathToLLVM.cpp
index 23e957288eb95e..668f8385ac2dcf 100644
--- a/mlir/lib/Conversion/MathToLLVM/MathToLLVM.cpp
+++ b/mlir/lib/Conversion/MathToLLVM/MathToLLVM.cpp
@@ -298,9 +298,9 @@ struct ConvertMathToLLVMPass
};
} // namespace
-void mlir::populateMathToLLVMConversionPatterns(LLVMTypeConverter &converter,
- RewritePatternSet &patterns,
- bool approximateLog1p) {
+void mlir::populateMathToLLVMConversionPatterns(
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns,
+ bool approximateLog1p) {
if (approximateLog1p)
patterns.add<Log1pOpLowering>(converter);
// clang-format off
diff --git a/mlir/lib/Conversion/MathToROCDL/MathToROCDL.cpp b/mlir/lib/Conversion/MathToROCDL/MathToROCDL.cpp
index 8330713ea66e5c..c17bfe4f71a98d 100644
--- a/mlir/lib/Conversion/MathToROCDL/MathToROCDL.cpp
+++ b/mlir/lib/Conversion/MathToROCDL/MathToROCDL.cpp
@@ -36,7 +36,7 @@ using namespace mlir;
#define DBGS() (llvm::dbgs() << "[" DEBUG_TYPE "]: ")
template <typename OpTy>
-static void populateOpPatterns(LLVMTypeConverter &converter,
+static void populateOpPatterns(const LLVMTypeConverter &converter,
RewritePatternSet &patterns, StringRef f32Func,
StringRef f64Func, StringRef f16Func,
StringRef f32ApproxFunc = "") {
@@ -45,8 +45,8 @@ static void populateOpPatterns(LLVMTypeConverter &converter,
f32ApproxFunc, f16Func);
}
-void mlir::populateMathToROCDLConversionPatterns(LLVMTypeConverter &converter,
- RewritePatternSet &patterns) {
+void mlir::populateMathToROCDLConversionPatterns(
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns) {
// Handled by mathToLLVM: math::AbsIOp
// Handled by mathToLLVM: math::AbsFOp
// Handled by mathToLLVM: math::CopySignOp
diff --git a/mlir/lib/Conversion/MathToSPIRV/MathToSPIRV.cpp b/mlir/lib/Conversion/MathToSPIRV/MathToSPIRV.cpp
index 52ff138bedf65b..6f948e80d5af8f 100644
--- a/mlir/lib/Conversion/MathToSPIRV/MathToSPIRV.cpp
+++ b/mlir/lib/Conversion/MathToSPIRV/MathToSPIRV.cpp
@@ -462,7 +462,7 @@ struct RoundOpPattern final : public OpConversionPattern<math::RoundOp> {
//===----------------------------------------------------------------------===//
namespace mlir {
-void populateMathToSPIRVPatterns(SPIRVTypeConverter &typeConverter,
+void populateMathToSPIRVPatterns(const SPIRVTypeConverter &typeConverter,
RewritePatternSet &patterns) {
// Core patterns
patterns.add<CopySignPattern>(typeConverter, patterns.getContext());
diff --git a/mlir/lib/Conversion/MemRefToEmitC/MemRefToEmitC.cpp b/mlir/lib/Conversion/MemRefToEmitC/MemRefToEmitC.cpp
index 1392bf924002ee..2b7ac4b529cf0d 100644
--- a/mlir/lib/Conversion/MemRefToEmitC/MemRefToEmitC.cpp
+++ b/mlir/lib/Conversion/MemRefToEmitC/MemRefToEmitC.cpp
@@ -179,8 +179,8 @@ void mlir::populateMemRefToEmitCTypeConversion(TypeConverter &typeConverter) {
});
}
-void mlir::populateMemRefToEmitCConversionPatterns(RewritePatternSet &patterns,
- TypeConverter &converter) {
+void mlir::populateMemRefToEmitCConversionPatterns(
+ RewritePatternSet &patterns, const TypeConverter &converter) {
patterns.add<ConvertAlloca, ConvertGlobal, ConvertGetGlobal, ConvertLoad,
ConvertStore>(converter, patterns.getContext());
}
diff --git a/mlir/lib/Conversion/MemRefToLLVM/MemRefToLLVM.cpp b/mlir/lib/Conversion/MemRefToLLVM/MemRefToLLVM.cpp
index 054827d40f0f32..4bfa536cc8a44a 100644
--- a/mlir/lib/Conversion/MemRefToLLVM/MemRefToLLVM.cpp
+++ b/mlir/lib/Conversion/MemRefToLLVM/MemRefToLLVM.cpp
@@ -1667,7 +1667,7 @@ class ExtractStridedMetadataOpLowering
} // namespace
void mlir::populateFinalizeMemRefToLLVMConversionPatterns(
- LLVMTypeConverter &converter, RewritePatternSet &patterns) {
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns) {
// clang-format off
patterns.add<
AllocaOpLowering,
diff --git a/mlir/lib/Conversion/MemRefToSPIRV/MemRefToSPIRV.cpp b/mlir/lib/Conversion/MemRefToSPIRV/MemRefToSPIRV.cpp
index 90b0d727ddee7f..285398311fd197 100644
--- a/mlir/lib/Conversion/MemRefToSPIRV/MemRefToSPIRV.cpp
+++ b/mlir/lib/Conversion/MemRefToSPIRV/MemRefToSPIRV.cpp
@@ -926,7 +926,7 @@ LogicalResult ReinterpretCastPattern::matchAndRewrite(
//===----------------------------------------------------------------------===//
namespace mlir {
-void populateMemRefToSPIRVPatterns(SPIRVTypeConverter &typeConverter,
+void populateMemRefToSPIRVPatterns(const SPIRVTypeConverter &typeConverter,
RewritePatternSet &patterns) {
patterns.add<AllocaOpPattern, AllocOpPattern, AtomicRMWOpPattern,
DeallocOpPattern, IntLoadOpPattern, IntStoreOpPattern,
diff --git a/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp b/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp
index ab47c9ae8fe79b..34a6b1d506540d 100644
--- a/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp
+++ b/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp
@@ -1701,8 +1701,8 @@ struct NVGPURcpOpLowering : public ConvertOpToLLVMPattern<nvgpu::RcpOp> {
};
} // namespace
-void mlir::populateNVGPUToNVVMConversionPatterns(LLVMTypeConverter &converter,
- RewritePatternSet &patterns) {
+void mlir::populateNVGPUToNVVMConversionPatterns(
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns) {
patterns.add<
NVGPUMBarrierCreateLowering, // nvgpu.mbarrier.create
NVGPUMBarrierInitLowering, // nvgpu.mbarrier.init
diff --git a/mlir/lib/Conversion/OpenMPToLLVM/OpenMPToLLVM.cpp b/mlir/lib/Conversion/OpenMPToLLVM/OpenMPToLLVM.cpp
index d6b4ec8584b082..bf1b59a9913f49 100644
--- a/mlir/lib/Conversion/OpenMPToLLVM/OpenMPToLLVM.cpp
+++ b/mlir/lib/Conversion/OpenMPToLLVM/OpenMPToLLVM.cpp
@@ -221,7 +221,7 @@ void MultiRegionOpConversion<omp::PrivateClauseOp>::forwardOpAttrs(
} // namespace
void mlir::configureOpenMPToLLVMConversionLegality(
- ConversionTarget &target, LLVMTypeConverter &typeConverter) {
+ ConversionTarget &target, const LLVMTypeConverter &typeConverter) {
target.addDynamicallyLegalOp<
omp::AtomicReadOp, omp::AtomicWriteOp, omp::CancellationPointOp,
omp::CancelOp, omp::CriticalDeclareOp, omp::FlushOp, omp::MapBoundsOp,
diff --git a/mlir/lib/Conversion/SCFToSPIRV/SCFToSPIRV.cpp b/mlir/lib/Conversion/SCFToSPIRV/SCFToSPIRV.cpp
index f62de1f17a6668..31d8cd2206148f 100644
--- a/mlir/lib/Conversion/SCFToSPIRV/SCFToSPIRV.cpp
+++ b/mlir/lib/Conversion/SCFToSPIRV/SCFToSPIRV.cpp
@@ -96,7 +96,7 @@ Region::iterator getBlockIt(Region ®ion, unsigned index) {
template <typename OpTy>
class SCFToSPIRVPattern : public OpConversionPattern<OpTy> {
public:
- SCFToSPIRVPattern(MLIRContext *context, SPIRVTypeConverter &converter,
+ SCFToSPIRVPattern(MLIRContext *context, const SPIRVTypeConverter &converter,
ScfToSPIRVContextImpl *scfToSPIRVContext)
: OpConversionPattern<OpTy>::OpConversionPattern(converter, context),
scfToSPIRVContext(scfToSPIRVContext), typeConverter(converter) {}
@@ -117,7 +117,7 @@ class SCFToSPIRVPattern : public OpConversionPattern<OpTy> {
// conversion. There isn't a straightforward way to do that yet, as when
// converting types, ops aren't taken into consideration. Therefore, we just
// bypass the framework's type conversion for now.
- SPIRVTypeConverter &typeConverter;
+ const SPIRVTypeConverter &typeConverter;
};
//===----------------------------------------------------------------------===//
@@ -436,7 +436,7 @@ struct WhileOpConversion final : SCFToSPIRVPattern<scf::WhileOp> {
// Public API
//===----------------------------------------------------------------------===//
-void mlir::populateSCFToSPIRVPatterns(SPIRVTypeConverter &typeConverter,
+void mlir::populateSCFToSPIRVPatterns(const SPIRVTypeConverter &typeConverter,
ScfToSPIRVContext &scfToSPIRVContext,
RewritePatternSet &patterns) {
patterns.add<ForOpConversion, IfOpConversion, TerminatorOpConversion,
diff --git a/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVM.cpp b/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVM.cpp
index 6ae607f75adbd5..e36e3951a31ecc 100644
--- a/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVM.cpp
+++ b/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVM.cpp
@@ -149,7 +149,7 @@ static Value optionallyTruncateOrExtend(Location loc, Value value,
/// Broadcasts the value to vector with `numElements` number of elements.
static Value broadcast(Location loc, Value toBroadcast, unsigned numElements,
- LLVMTypeConverter &typeConverter,
+ const LLVMTypeConverter &typeConverter,
ConversionPatternRewriter &rewriter) {
auto vectorType = VectorType::get(numElements, toBroadcast.getType());
auto llvmVectorType = typeConverter.convertType(vectorType);
@@ -166,7 +166,7 @@ static Value broadcast(Location loc, Value toBroadcast, unsigned numElements,
/// Broadcasts the value. If `srcType` is a scalar, the value remains unchanged.
static Value optionallyBroadcast(Location loc, Value value, Type srcType,
- LLVMTypeConverter &typeConverter,
+ const LLVMTypeConverter &typeConverter,
ConversionPatternRewriter &rewriter) {
if (auto vectorType = dyn_cast<VectorType>(srcType)) {
unsigned numElements = vectorType.getNumElements();
@@ -186,7 +186,8 @@ static Value optionallyBroadcast(Location loc, Value value, Type srcType,
/// Then cast `Offset` and `Count` if their bit width is different
/// from `Base` bit width.
static Value processCountOrOffset(Location loc, Value value, Type srcType,
- Type dstType, LLVMTypeConverter &converter,
+ Type dstType,
+ const LLVMTypeConverter &converter,
ConversionPatternRewriter &rewriter) {
Value broadcasted =
optionallyBroadcast(loc, value, srcType, converter, rewriter);
@@ -196,7 +197,7 @@ static Value processCountOrOffset(Location loc, Value value, Type srcType,
/// Converts SPIR-V struct with a regular (according to `VulkanLayoutUtils`)
/// offset to LLVM struct. Otherwise, the conversion is not supported.
static Type convertStructTypeWithOffset(spirv::StructType type,
- LLVMTypeConverter &converter) {
+ const LLVMTypeConverter &converter) {
if (type != VulkanLayoutUtils::decorateType(type))
return nullptr;
@@ -209,7 +210,7 @@ static Type convertStructTypeWithOffset(spirv::StructType type,
/// Converts SPIR-V struct with no offset to packed LLVM struct.
static Type convertStructTypePacked(spirv::StructType type,
- LLVMTypeConverter &converter) {
+ const LLVMTypeConverter &converter) {
SmallVector<Type> elementsVector;
if (failed(converter.convertTypes(type.getElementTypes(), elementsVector)))
return nullptr;
@@ -226,11 +227,10 @@ static Value createI32ConstantOf(Location loc, PatternRewriter &rewriter,
}
/// Utility for `spirv.Load` and `spirv.Store` conversion.
-static LogicalResult replaceWithLoadOrStore(Operation *op, ValueRange operands,
- ConversionPatternRewriter &rewriter,
- LLVMTypeConverter &typeConverter,
- unsigned alignment, bool isVolatile,
- bool isNonTemporal) {
+static LogicalResult replaceWithLoadOrStore(
+ Operation *op, ValueRange operands, ConversionPatternRewriter &rewriter,
+ const LLVMTypeConverter &typeConverter, unsigned alignment, bool isVolatile,
+ bool isNonTemporal) {
if (auto loadOp = dyn_cast<spirv::LoadOp>(op)) {
auto dstType = typeConverter.convertType(loadOp.getType());
if (!dstType)
@@ -271,7 +271,7 @@ static std::optional<Type> convertArrayType(spirv::ArrayType type,
/// Converts SPIR-V pointer type to LLVM pointer. Pointer's storage class is not
/// modelled at the moment.
static Type convertPointerType(spirv::PointerType type,
- LLVMTypeConverter &converter,
+ const LLVMTypeConverter &converter,
spirv::ClientAPI clientAPI) {
unsigned addressSpace =
storageClassToAddressSpace(clientAPI, type.getStorageClass());
@@ -292,7 +292,7 @@ static std::optional<Type> convertRuntimeArrayType(spirv::RuntimeArrayType type,
/// Converts SPIR-V struct to LLVM struct. There is no support of structs with
/// member decorations. Also, only natural offset is supported.
static Type convertStructType(spirv::StructType type,
- LLVMTypeConverter &converter) {
+ const LLVMTypeConverter &converter) {
SmallVector<spirv::StructType::MemberDecorationInfo, 4> memberDecorations;
type.getMemberDecorations(memberDecorations);
if (!memberDecorations.empty())
@@ -315,20 +315,21 @@ class AccessChainPattern : public SPIRVToLLVMConversion<spirv::AccessChainOp> {
LogicalResult
matchAndRewrite(spirv::AccessChainOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- auto dstType = typeConverter.convertType(op.getComponentPtr().getType());
+ auto dstType =
+ getTypeConverter()->convertType(op.getComponentPtr().getType());
if (!dstType)
return rewriter.notifyMatchFailure(op, "type conversion failed");
// To use GEP we need to add a first 0 index to go through the pointer.
auto indices = llvm::to_vector<4>(adaptor.getIndices());
Type indexType = op.getIndices().front().getType();
- auto llvmIndexType = typeConverter.convertType(indexType);
+ auto llvmIndexType = getTypeConverter()->convertType(indexType);
if (!llvmIndexType)
return rewriter.notifyMatchFailure(op, "type conversion failed");
Value zero = rewriter.create<LLVM::ConstantOp>(
op.getLoc(), llvmIndexType, rewriter.getIntegerAttr(indexType, 0));
indices.insert(indices.begin(), zero);
- auto elementType = typeConverter.convertType(
+ auto elementType = getTypeConverter()->convertType(
cast<spirv::PointerType>(op.getBasePtr().getType()).getPointeeType());
if (!elementType)
return rewriter.notifyMatchFailure(op, "type conversion failed");
@@ -345,7 +346,7 @@ class AddressOfPattern : public SPIRVToLLVMConversion<spirv::AddressOfOp> {
LogicalResult
matchAndRewrite(spirv::AddressOfOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- auto dstType = typeConverter.convertType(op.getPointer().getType());
+ auto dstType = getTypeConverter()->convertType(op.getPointer().getType());
if (!dstType)
return rewriter.notifyMatchFailure(op, "type conversion failed");
rewriter.replaceOpWithNewOp<LLVM::AddressOfOp>(op, dstType,
@@ -363,16 +364,16 @@ class BitFieldInsertPattern
matchAndRewrite(spirv::BitFieldInsertOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
auto srcType = op.getType();
- auto dstType = typeConverter.convertType(srcType);
+ auto dstType = getTypeConverter()->convertType(srcType);
if (!dstType)
return rewriter.notifyMatchFailure(op, "type conversion failed");
Location loc = op.getLoc();
// Process `Offset` and `Count`: broadcast and extend/truncate if needed.
Value offset = processCountOrOffset(loc, op.getOffset(), srcType, dstType,
- typeConverter, rewriter);
+ *getTypeConverter(), rewriter);
Value count = processCountOrOffset(loc, op.getCount(), srcType, dstType,
- typeConverter, rewriter);
+ *getTypeConverter(), rewriter);
// Create a mask with bits set outside [Offset, Offset + Count - 1].
Value minusOne = createConstantAllBitsSet(loc, srcType, dstType, rewriter);
@@ -410,7 +411,7 @@ class ConstantScalarAndVectorPattern
if (!isa<VectorType>(srcType) && !srcType.isIntOrFloat())
return failure();
- auto dstType = typeConverter.convertType(srcType);
+ auto dstType = getTypeConverter()->convertType(srcType);
if (!dstType)
return rewriter.notifyMatchFailure(constOp, "type conversion failed");
@@ -451,16 +452,16 @@ class BitFieldSExtractPattern
matchAndRewrite(spirv::BitFieldSExtractOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
auto srcType = op.getType();
- auto dstType = typeConverter.convertType(srcType);
+ auto dstType = getTypeConverter()->convertType(srcType);
if (!dstType)
return rewriter.notifyMatchFailure(op, "type conversion failed");
Location loc = op.getLoc();
// Process `Offset` and `Count`: broadcast and extend/truncate if needed.
Value offset = processCountOrOffset(loc, op.getOffset(), srcType, dstType,
- typeConverter, rewriter);
+ *getTypeConverter(), rewriter);
Value count = processCountOrOffset(loc, op.getCount(), srcType, dstType,
- typeConverter, rewriter);
+ *getTypeConverter(), rewriter);
// Create a constant that holds the size of the `Base`.
IntegerType integerType;
@@ -504,16 +505,16 @@ class BitFieldUExtractPattern
matchAndRewrite(spirv::BitFieldUExtractOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
auto srcType = op.getType();
- auto dstType = typeConverter.convertType(srcType);
+ auto dstType = getTypeConverter()->convertType(srcType);
if (!dstType)
return rewriter.notifyMatchFailure(op, "type conversion failed");
Location loc = op.getLoc();
// Process `Offset` and `Count`: broadcast and extend/truncate if needed.
Value offset = processCountOrOffset(loc, op.getOffset(), srcType, dstType,
- typeConverter, rewriter);
+ *getTypeConverter(), rewriter);
Value count = processCountOrOffset(loc, op.getCount(), srcType, dstType,
- typeConverter, rewriter);
+ *getTypeConverter(), rewriter);
// Create a mask with bits set at [0, Count - 1].
Value minusOne = createConstantAllBitsSet(loc, srcType, dstType, rewriter);
@@ -580,7 +581,7 @@ class CompositeExtractPattern
LogicalResult
matchAndRewrite(spirv::CompositeExtractOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- auto dstType = this->typeConverter.convertType(op.getType());
+ auto dstType = this->getTypeConverter()->convertType(op.getType());
if (!dstType)
return rewriter.notifyMatchFailure(op, "type conversion failed");
@@ -612,7 +613,7 @@ class CompositeInsertPattern
LogicalResult
matchAndRewrite(spirv::CompositeInsertOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- auto dstType = this->typeConverter.convertType(op.getType());
+ auto dstType = this->getTypeConverter()->convertType(op.getType());
if (!dstType)
return rewriter.notifyMatchFailure(op, "type conversion failed");
@@ -643,7 +644,7 @@ class DirectConversionPattern : public SPIRVToLLVMConversion<SPIRVOp> {
LogicalResult
matchAndRewrite(SPIRVOp op, typename SPIRVOp::Adaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- auto dstType = this->typeConverter.convertType(op.getType());
+ auto dstType = this->getTypeConverter()->convertType(op.getType());
if (!dstType)
return rewriter.notifyMatchFailure(op, "type conversion failed");
rewriter.template replaceOpWithNewOp<LLVMOp>(
@@ -749,7 +750,7 @@ class GlobalVariablePattern
return failure();
auto srcType = cast<spirv::PointerType>(op.getType());
- auto dstType = typeConverter.convertType(srcType.getPointeeType());
+ auto dstType = getTypeConverter()->convertType(srcType.getPointeeType());
if (!dstType)
return rewriter.notifyMatchFailure(op, "type conversion failed");
@@ -810,7 +811,7 @@ class IndirectCastPattern : public SPIRVToLLVMConversion<SPIRVOp> {
Type fromType = op.getOperand().getType();
Type toType = op.getType();
- auto dstType = this->typeConverter.convertType(toType);
+ auto dstType = this->getTypeConverter()->convertType(toType);
if (!dstType)
return rewriter.notifyMatchFailure(op, "type conversion failed");
@@ -846,7 +847,7 @@ class FunctionCallPattern
}
// Function returns a single result.
- auto dstType = typeConverter.convertType(callOp.getType(0));
+ auto dstType = getTypeConverter()->convertType(callOp.getType(0));
if (!dstType)
return rewriter.notifyMatchFailure(callOp, "type conversion failed");
auto newOp = rewriter.replaceOpWithNewOp<LLVM::CallOp>(
@@ -868,7 +869,7 @@ class FComparePattern : public SPIRVToLLVMConversion<SPIRVOp> {
matchAndRewrite(SPIRVOp op, typename SPIRVOp::Adaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- auto dstType = this->typeConverter.convertType(op.getType());
+ auto dstType = this->getTypeConverter()->convertType(op.getType());
if (!dstType)
return rewriter.notifyMatchFailure(op, "type conversion failed");
@@ -888,7 +889,7 @@ class IComparePattern : public SPIRVToLLVMConversion<SPIRVOp> {
matchAndRewrite(SPIRVOp op, typename SPIRVOp::Adaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- auto dstType = this->typeConverter.convertType(op.getType());
+ auto dstType = this->getTypeConverter()->convertType(op.getType());
if (!dstType)
return rewriter.notifyMatchFailure(op, "type conversion failed");
@@ -907,7 +908,7 @@ class InverseSqrtPattern
matchAndRewrite(spirv::GLInverseSqrtOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
auto srcType = op.getType();
- auto dstType = typeConverter.convertType(srcType);
+ auto dstType = getTypeConverter()->convertType(srcType);
if (!dstType)
return rewriter.notifyMatchFailure(op, "type conversion failed");
@@ -930,7 +931,7 @@ class LoadStorePattern : public SPIRVToLLVMConversion<SPIRVOp> {
ConversionPatternRewriter &rewriter) const override {
if (!op.getMemoryAccess()) {
return replaceWithLoadOrStore(op, adaptor.getOperands(), rewriter,
- this->typeConverter, /*alignment=*/0,
+ *this->getTypeConverter(), /*alignment=*/0,
/*isVolatile=*/false,
/*isNonTemporal=*/false);
}
@@ -945,8 +946,8 @@ class LoadStorePattern : public SPIRVToLLVMConversion<SPIRVOp> {
bool isNonTemporal = memoryAccess == spirv::MemoryAccess::Nontemporal;
bool isVolatile = memoryAccess == spirv::MemoryAccess::Volatile;
return replaceWithLoadOrStore(op, adaptor.getOperands(), rewriter,
- this->typeConverter, alignment, isVolatile,
- isNonTemporal);
+ *this->getTypeConverter(), alignment,
+ isVolatile, isNonTemporal);
}
default:
// There is no support of other memory access attributes.
@@ -965,7 +966,7 @@ class NotPattern : public SPIRVToLLVMConversion<SPIRVOp> {
matchAndRewrite(SPIRVOp notOp, typename SPIRVOp::Adaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
auto srcType = notOp.getType();
- auto dstType = this->typeConverter.convertType(srcType);
+ auto dstType = this->getTypeConverter()->convertType(srcType);
if (!dstType)
return rewriter.notifyMatchFailure(notOp, "type conversion failed");
@@ -1196,7 +1197,7 @@ class ShiftPattern : public SPIRVToLLVMConversion<SPIRVOp> {
matchAndRewrite(SPIRVOp op, typename SPIRVOp::Adaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- auto dstType = this->typeConverter.convertType(op.getType());
+ auto dstType = this->getTypeConverter()->convertType(op.getType());
if (!dstType)
return rewriter.notifyMatchFailure(op, "type conversion failed");
@@ -1247,7 +1248,7 @@ class TanPattern : public SPIRVToLLVMConversion<spirv::GLTanOp> {
LogicalResult
matchAndRewrite(spirv::GLTanOp tanOp, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- auto dstType = typeConverter.convertType(tanOp.getType());
+ auto dstType = getTypeConverter()->convertType(tanOp.getType());
if (!dstType)
return rewriter.notifyMatchFailure(tanOp, "type conversion failed");
@@ -1273,7 +1274,7 @@ class TanhPattern : public SPIRVToLLVMConversion<spirv::GLTanhOp> {
matchAndRewrite(spirv::GLTanhOp tanhOp, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
auto srcType = tanhOp.getType();
- auto dstType = typeConverter.convertType(srcType);
+ auto dstType = getTypeConverter()->convertType(srcType);
if (!dstType)
return rewriter.notifyMatchFailure(tanhOp, "type conversion failed");
@@ -1307,21 +1308,21 @@ class VariablePattern : public SPIRVToLLVMConversion<spirv::VariableOp> {
if (init && !pointerTo.isIntOrFloat() && !isa<VectorType>(pointerTo))
return failure();
- auto dstType = typeConverter.convertType(srcType);
+ auto dstType = getTypeConverter()->convertType(srcType);
if (!dstType)
return rewriter.notifyMatchFailure(varOp, "type conversion failed");
Location loc = varOp.getLoc();
Value size = createI32ConstantOf(loc, rewriter, 1);
if (!init) {
- auto elementType = typeConverter.convertType(pointerTo);
+ auto elementType = getTypeConverter()->convertType(pointerTo);
if (!elementType)
return rewriter.notifyMatchFailure(varOp, "type conversion failed");
rewriter.replaceOpWithNewOp<LLVM::AllocaOp>(varOp, dstType, elementType,
size);
return success();
}
- auto elementType = typeConverter.convertType(pointerTo);
+ auto elementType = getTypeConverter()->convertType(pointerTo);
if (!elementType)
return rewriter.notifyMatchFailure(varOp, "type conversion failed");
Value allocated =
@@ -1344,7 +1345,7 @@ class BitcastConversionPattern
LogicalResult
matchAndRewrite(spirv::BitcastOp bitcastOp, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- auto dstType = typeConverter.convertType(bitcastOp.getType());
+ auto dstType = getTypeConverter()->convertType(bitcastOp.getType());
if (!dstType)
return rewriter.notifyMatchFailure(bitcastOp, "type conversion failed");
@@ -1377,7 +1378,7 @@ class FuncConversionPattern : public SPIRVToLLVMConversion<spirv::FuncOp> {
auto funcType = funcOp.getFunctionType();
TypeConverter::SignatureConversion signatureConverter(
funcType.getNumInputs());
- auto llvmType = typeConverter.convertFunctionSignature(
+ auto llvmType = getTypeConverter()->convertFunctionSignature(
funcType, /*isVariadic=*/false, /*useBarePtrCallConv=*/false,
signatureConverter);
if (!llvmType)
@@ -1418,8 +1419,8 @@ class FuncConversionPattern : public SPIRVToLLVMConversion<spirv::FuncOp> {
rewriter.inlineRegionBefore(funcOp.getBody(), newFuncOp.getBody(),
newFuncOp.end());
- if (failed(rewriter.convertRegionTypes(&newFuncOp.getBody(), typeConverter,
- &signatureConverter))) {
+ if (failed(rewriter.convertRegionTypes(
+ &newFuncOp.getBody(), *getTypeConverter(), &signatureConverter))) {
return failure();
}
rewriter.eraseOp(funcOp);
@@ -1474,7 +1475,7 @@ class VectorShufflePattern
return success();
}
- auto dstType = typeConverter.convertType(op.getType());
+ auto dstType = getTypeConverter()->convertType(op.getType());
if (!dstType)
return rewriter.notifyMatchFailure(op, "type conversion failed");
auto scalarType = cast<VectorType>(dstType).getElementType();
@@ -1535,7 +1536,7 @@ void mlir::populateSPIRVToLLVMTypeConversion(LLVMTypeConverter &typeConverter,
}
void mlir::populateSPIRVToLLVMConversionPatterns(
- LLVMTypeConverter &typeConverter, RewritePatternSet &patterns,
+ const LLVMTypeConverter &typeConverter, RewritePatternSet &patterns,
spirv::ClientAPI clientAPI) {
patterns.add<
// Arithmetic ops
@@ -1653,12 +1654,12 @@ void mlir::populateSPIRVToLLVMConversionPatterns(
}
void mlir::populateSPIRVToLLVMFunctionConversionPatterns(
- LLVMTypeConverter &typeConverter, RewritePatternSet &patterns) {
+ const LLVMTypeConverter &typeConverter, RewritePatternSet &patterns) {
patterns.add<FuncConversionPattern>(patterns.getContext(), typeConverter);
}
void mlir::populateSPIRVToLLVMModuleConversionPatterns(
- LLVMTypeConverter &typeConverter, RewritePatternSet &patterns) {
+ const LLVMTypeConverter &typeConverter, RewritePatternSet &patterns) {
patterns.add<ModuleConversionPattern>(patterns.getContext(), typeConverter);
}
diff --git a/mlir/lib/Conversion/TensorToSPIRV/TensorToSPIRV.cpp b/mlir/lib/Conversion/TensorToSPIRV/TensorToSPIRV.cpp
index 468fffdd2df91b..ef3f4bc22fb898 100644
--- a/mlir/lib/Conversion/TensorToSPIRV/TensorToSPIRV.cpp
+++ b/mlir/lib/Conversion/TensorToSPIRV/TensorToSPIRV.cpp
@@ -35,7 +35,7 @@ namespace {
class TensorExtractPattern final
: public OpConversionPattern<tensor::ExtractOp> {
public:
- TensorExtractPattern(TypeConverter &typeConverter, MLIRContext *context,
+ TensorExtractPattern(const TypeConverter &typeConverter, MLIRContext *context,
int64_t threshold, PatternBenefit benefit = 1)
: OpConversionPattern(typeConverter, context, benefit),
byteCountThreshold(threshold) {}
@@ -103,9 +103,9 @@ class TensorExtractPattern final
// Pattern population
//===----------------------------------------------------------------------===//
-void mlir::populateTensorToSPIRVPatterns(SPIRVTypeConverter &typeConverter,
- int64_t byteCountThreshold,
- RewritePatternSet &patterns) {
+void mlir::populateTensorToSPIRVPatterns(
+ const SPIRVTypeConverter &typeConverter, int64_t byteCountThreshold,
+ RewritePatternSet &patterns) {
patterns.add<TensorExtractPattern>(typeConverter, patterns.getContext(),
byteCountThreshold);
}
diff --git a/mlir/lib/Conversion/TosaToLinalg/TosaToLinalg.cpp b/mlir/lib/Conversion/TosaToLinalg/TosaToLinalg.cpp
index 01fdd57260797b..c88f4db27c304e 100644
--- a/mlir/lib/Conversion/TosaToLinalg/TosaToLinalg.cpp
+++ b/mlir/lib/Conversion/TosaToLinalg/TosaToLinalg.cpp
@@ -2588,7 +2588,7 @@ struct FFT2dConverter final : OpRewritePattern<FFT2dOp> {
} // namespace
void mlir::tosa::populateTosaToLinalgConversionPatterns(
- TypeConverter &converter, RewritePatternSet *patterns) {
+ const TypeConverter &converter, RewritePatternSet *patterns) {
// We have multiple resize coverters to handle degenerate cases.
patterns->add<GenericResizeConverter>(patterns->getContext(),
diff --git a/mlir/lib/Conversion/TosaToTensor/TosaToTensor.cpp b/mlir/lib/Conversion/TosaToTensor/TosaToTensor.cpp
index 3f104ed1e3f7fb..2f2838751d70f0 100644
--- a/mlir/lib/Conversion/TosaToTensor/TosaToTensor.cpp
+++ b/mlir/lib/Conversion/TosaToTensor/TosaToTensor.cpp
@@ -35,10 +35,10 @@ TensorType inferReshapeInputType(TypedValue<TensorType> input,
return input.getType();
// The input type must be cast into a tensor with the same rank and all static
- // dimensions set to 1. This prevents the generation of a tensor.collapse_shape
- // op that converts a dynamically shaped tensor into a 0D tensor. While such
- // construct is not incorrect on its own, bufferization cannot properly handle
- // it at the moment, so we avoid it.
+ // dimensions set to 1. This prevents the generation of a
+ // tensor.collapse_shape op that converts a dynamically shaped tensor into a
+ // 0D tensor. While such construct is not incorrect on its own, bufferization
+ // cannot properly handle it at the moment, so we avoid it.
SmallVector<int64_t> shape(input.getType().getRank(), 1);
return input.getType().clone(shape);
}
@@ -57,29 +57,31 @@ TensorType inferReshapeExpandedType(TensorType inputType,
int64_t totalSize = inputIsStatic ? inputType.getNumElements() : -1;
// Compute result shape
- auto resultShape = llvm::map_to_vector(newShape, [&](int64_t size) -> int64_t {
- // If this is not a placeholder, do not change it
- if (size >= 0)
- return size;
-
- // If we do not know the total size of the tensor, keep this dimension
- // dynamic in the result shape.
- if (!inputIsStatic)
- return ShapedType::kDynamic;
-
- // Calculate the product of all elements in 'newShape' except for the -1
- // placeholder, which we discard by negating the result.
- int64_t totalSizeNoPlaceholder = -std::accumulate(
- newShape.begin(), newShape.end(), 1, std::multiplies<int64_t>());
-
- // If there is a 0 component in 'newShape', resolve the placeholder as 0.
- if (totalSizeNoPlaceholder == 0)
- return 0;
-
- // Resolve the placeholder as the quotient between the total tensor size and
- // the product of all other sizes.
- return totalSize / totalSizeNoPlaceholder;
- });
+ auto resultShape =
+ llvm::map_to_vector(newShape, [&](int64_t size) -> int64_t {
+ // If this is not a placeholder, do not change it
+ if (size >= 0)
+ return size;
+
+ // If we do not know the total size of the tensor, keep this dimension
+ // dynamic in the result shape.
+ if (!inputIsStatic)
+ return ShapedType::kDynamic;
+
+ // Calculate the product of all elements in 'newShape' except for the -1
+ // placeholder, which we discard by negating the result.
+ int64_t totalSizeNoPlaceholder = -std::accumulate(
+ newShape.begin(), newShape.end(), 1, std::multiplies<int64_t>());
+
+ // If there is a 0 component in 'newShape', resolve the placeholder as
+ // 0.
+ if (totalSizeNoPlaceholder == 0)
+ return 0;
+
+ // Resolve the placeholder as the quotient between the total tensor size
+ // and the product of all other sizes.
+ return totalSize / totalSizeNoPlaceholder;
+ });
bool resultIsStatic = !ShapedType::isDynamicShape(resultShape);
@@ -107,7 +109,8 @@ TensorType inferReshapeCollapsedType(TensorType lhsType, TensorType rhsType) {
if (lhsShape.empty() || rhsShape.empty())
return lhsType.clone(ArrayRef<int64_t>{});
- if (ShapedType::isDynamicShape(lhsShape) || ShapedType::isDynamicShape(rhsShape))
+ if (ShapedType::isDynamicShape(lhsShape) ||
+ ShapedType::isDynamicShape(rhsShape))
return lhsType.clone({ShapedType::kDynamic});
SmallVector<int64_t> intermediateShape;
@@ -149,14 +152,16 @@ TensorType inferReshapeCollapsedType(TensorType lhsType, TensorType rhsType) {
}
SmallVector<ReassociationExprs>
-createReassociationMapForCollapse(OpBuilder &builder, Type srcType, Type dstType) {
+createReassociationMapForCollapse(OpBuilder &builder, Type srcType,
+ Type dstType) {
auto srcShape = cast<TensorType>(srcType).getShape();
auto dstShape = cast<TensorType>(dstType).getShape();
if (srcShape.empty() || dstShape.empty())
return {};
- if (ShapedType::isDynamicShape(srcShape) || ShapedType::isDynamicShape(dstShape)) {
+ if (ShapedType::isDynamicShape(srcShape) ||
+ ShapedType::isDynamicShape(dstShape)) {
assert(dstShape.size() == 1);
SmallVector<AffineExpr, 2> exprs;
for (auto i : llvm::seq<int64_t>(srcShape.size()))
@@ -243,14 +248,16 @@ class ReshapeConverter : public OpConversionPattern<tosa::ReshapeOp> {
auto collapsedType = inferReshapeCollapsedType(inputType, expandedType);
// Cast input if needed
- auto castInput = rewriter.createOrFold<tensor::CastOp>(loc, inputType, input);
+ auto castInput =
+ rewriter.createOrFold<tensor::CastOp>(loc, inputType, input);
// Emit collaspe-expand pair
auto collapsed = createCollapse(rewriter, loc, collapsedType, castInput);
auto expanded = createExpand(rewriter, loc, expandedType, collapsed);
// Cast to final result type if needed
- auto result = rewriter.createOrFold<tensor::CastOp>(loc, resultType, expanded);
+ auto result =
+ rewriter.createOrFold<tensor::CastOp>(loc, resultType, expanded);
rewriter.replaceOp(reshape, result);
return success();
}
@@ -438,7 +445,7 @@ struct ConcatConverter : public OpConversionPattern<tosa::ConcatOp> {
} // namespace
void mlir::tosa::populateTosaToTensorConversionPatterns(
- TypeConverter &converter, RewritePatternSet *patterns) {
+ const TypeConverter &converter, RewritePatternSet *patterns) {
patterns
->add<ConcatConverter, PadConverter, ReshapeConverter, SliceConverter>(
converter, patterns->getContext());
diff --git a/mlir/lib/Conversion/UBToLLVM/UBToLLVM.cpp b/mlir/lib/Conversion/UBToLLVM/UBToLLVM.cpp
index 0051333a35dcdc..9921a06778dd7d 100644
--- a/mlir/lib/Conversion/UBToLLVM/UBToLLVM.cpp
+++ b/mlir/lib/Conversion/UBToLLVM/UBToLLVM.cpp
@@ -91,8 +91,8 @@ struct UBToLLVMConversionPass
// Pattern Population
//===----------------------------------------------------------------------===//
-void mlir::ub::populateUBToLLVMConversionPatterns(LLVMTypeConverter &converter,
- RewritePatternSet &patterns) {
+void mlir::ub::populateUBToLLVMConversionPatterns(
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns) {
patterns.add<PoisonOpLowering>(converter);
}
diff --git a/mlir/lib/Conversion/UBToSPIRV/UBToSPIRV.cpp b/mlir/lib/Conversion/UBToSPIRV/UBToSPIRV.cpp
index 001b7fefb175df..a3806189e40608 100644
--- a/mlir/lib/Conversion/UBToSPIRV/UBToSPIRV.cpp
+++ b/mlir/lib/Conversion/UBToSPIRV/UBToSPIRV.cpp
@@ -79,6 +79,6 @@ struct UBToSPIRVConversionPass final
//===----------------------------------------------------------------------===//
void mlir::ub::populateUBToSPIRVConversionPatterns(
- SPIRVTypeConverter &converter, RewritePatternSet &patterns) {
+ const SPIRVTypeConverter &converter, RewritePatternSet &patterns) {
patterns.add<PoisonOpLowering>(converter, patterns.getContext());
}
diff --git a/mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVM.cpp b/mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVM.cpp
index 687061e9988f8f..a150e2be11736a 100644
--- a/mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVM.cpp
+++ b/mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVM.cpp
@@ -1881,7 +1881,7 @@ struct VectorStepOpLowering : public ConvertOpToLLVMPattern<vector::StepOp> {
/// Populate the given list with patterns that convert from Vector to LLVM.
void mlir::populateVectorToLLVMConversionPatterns(
- LLVMTypeConverter &converter, RewritePatternSet &patterns,
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns,
bool reassociateFPReductions, bool force32BitVectorIndices) {
MLIRContext *ctx = converter.getDialect()->getContext();
patterns.add<VectorFMAOpNDRewritePattern>(ctx);
@@ -1909,7 +1909,7 @@ void mlir::populateVectorToLLVMConversionPatterns(
}
void mlir::populateVectorToLLVMMatrixConversionPatterns(
- LLVMTypeConverter &converter, RewritePatternSet &patterns) {
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns) {
patterns.add<VectorMatmulOpConversion>(converter);
patterns.add<VectorFlatTransposeOpConversion>(converter);
}
diff --git a/mlir/lib/Conversion/VectorToSPIRV/VectorToSPIRV.cpp b/mlir/lib/Conversion/VectorToSPIRV/VectorToSPIRV.cpp
index 21b8858989839b..6184225cb6285d 100644
--- a/mlir/lib/Conversion/VectorToSPIRV/VectorToSPIRV.cpp
+++ b/mlir/lib/Conversion/VectorToSPIRV/VectorToSPIRV.cpp
@@ -950,8 +950,8 @@ struct VectorStepOpConvert final : OpConversionPattern<vector::StepOp> {
#define CL_FLOAT_MAX_MIN_OPS spirv::CLFMaxOp, spirv::CLFMinOp
#define GL_FLOAT_MAX_MIN_OPS spirv::GLFMaxOp, spirv::GLFMinOp
-void mlir::populateVectorToSPIRVPatterns(SPIRVTypeConverter &typeConverter,
- RewritePatternSet &patterns) {
+void mlir::populateVectorToSPIRVPatterns(
+ const SPIRVTypeConverter &typeConverter, RewritePatternSet &patterns) {
patterns.add<
VectorBitcastConvert, VectorBroadcastConvert,
VectorExtractElementOpConvert, VectorExtractOpConvert,
diff --git a/mlir/lib/Dialect/AMX/Transforms/LegalizeForLLVMExport.cpp b/mlir/lib/Dialect/AMX/Transforms/LegalizeForLLVMExport.cpp
index a8b10f63315d41..c8cfcc3d945bec 100644
--- a/mlir/lib/Dialect/AMX/Transforms/LegalizeForLLVMExport.cpp
+++ b/mlir/lib/Dialect/AMX/Transforms/LegalizeForLLVMExport.cpp
@@ -203,7 +203,7 @@ struct TileMulIConversion : public ConvertOpToLLVMPattern<TileMulIOp> {
} // namespace
void mlir::populateAMXLegalizeForLLVMExportPatterns(
- LLVMTypeConverter &converter, RewritePatternSet &patterns) {
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns) {
patterns.add<TileZeroConversion, TileLoadConversion, TileStoreConversion,
TileMulFConversion, TileMulIConversion>(converter);
}
diff --git a/mlir/lib/Dialect/Arith/Transforms/EmulateNarrowType.cpp b/mlir/lib/Dialect/Arith/Transforms/EmulateNarrowType.cpp
index 1c1f65e580eb0c..4be0e06fe2a5e5 100644
--- a/mlir/lib/Dialect/Arith/Transforms/EmulateNarrowType.cpp
+++ b/mlir/lib/Dialect/Arith/Transforms/EmulateNarrowType.cpp
@@ -51,7 +51,8 @@ arith::NarrowTypeEmulationConverter::NarrowTypeEmulationConverter(
}
void arith::populateArithNarrowTypeEmulationPatterns(
- NarrowTypeEmulationConverter &typeConverter, RewritePatternSet &patterns) {
+ const NarrowTypeEmulationConverter &typeConverter,
+ RewritePatternSet &patterns) {
// Populate `func.*` conversion patterns.
populateFunctionOpInterfaceTypeConversionPattern<func::FuncOp>(patterns,
typeConverter);
diff --git a/mlir/lib/Dialect/Arith/Transforms/EmulateUnsupportedFloats.cpp b/mlir/lib/Dialect/Arith/Transforms/EmulateUnsupportedFloats.cpp
index b51444e884aaee..836ebb65e7d17b 100644
--- a/mlir/lib/Dialect/Arith/Transforms/EmulateUnsupportedFloats.cpp
+++ b/mlir/lib/Dialect/Arith/Transforms/EmulateUnsupportedFloats.cpp
@@ -41,7 +41,7 @@ struct EmulateUnsupportedFloatsPass
};
struct EmulateFloatPattern final : ConversionPattern {
- EmulateFloatPattern(TypeConverter &converter, MLIRContext *ctx)
+ EmulateFloatPattern(const TypeConverter &converter, MLIRContext *ctx)
: ConversionPattern(converter, Pattern::MatchAnyOpTypeTag(), 1, ctx) {}
LogicalResult match(Operation *op) const override;
@@ -106,12 +106,12 @@ void mlir::arith::populateEmulateUnsupportedFloatsConversions(
}
void mlir::arith::populateEmulateUnsupportedFloatsPatterns(
- RewritePatternSet &patterns, TypeConverter &converter) {
+ RewritePatternSet &patterns, const TypeConverter &converter) {
patterns.add<EmulateFloatPattern>(converter, patterns.getContext());
}
void mlir::arith::populateEmulateUnsupportedFloatsLegality(
- ConversionTarget &target, TypeConverter &converter) {
+ ConversionTarget &target, const TypeConverter &converter) {
// Don't try to legalize functions and other ops that don't need expansion.
target.markUnknownOpDynamicallyLegal([](Operation *op) { return true; });
target.addDynamicallyLegalDialect<arith::ArithDialect>(
diff --git a/mlir/lib/Dialect/Arith/Transforms/EmulateWideInt.cpp b/mlir/lib/Dialect/Arith/Transforms/EmulateWideInt.cpp
index 9c91dd71b3e15d..c48ded094eacff 100644
--- a/mlir/lib/Dialect/Arith/Transforms/EmulateWideInt.cpp
+++ b/mlir/lib/Dialect/Arith/Transforms/EmulateWideInt.cpp
@@ -1122,7 +1122,8 @@ arith::WideIntEmulationConverter::WideIntEmulationConverter(
}
void arith::populateArithWideIntEmulationPatterns(
- WideIntEmulationConverter &typeConverter, RewritePatternSet &patterns) {
+ const WideIntEmulationConverter &typeConverter,
+ RewritePatternSet &patterns) {
// Populate `func.*` conversion patterns.
populateFunctionOpInterfaceTypeConversionPattern<func::FuncOp>(patterns,
typeConverter);
diff --git a/mlir/lib/Dialect/ArmSVE/Transforms/LegalizeForLLVMExport.cpp b/mlir/lib/Dialect/ArmSVE/Transforms/LegalizeForLLVMExport.cpp
index 10f39a0855f5f5..845a32c4d97b5e 100644
--- a/mlir/lib/Dialect/ArmSVE/Transforms/LegalizeForLLVMExport.cpp
+++ b/mlir/lib/Dialect/ArmSVE/Transforms/LegalizeForLLVMExport.cpp
@@ -200,7 +200,7 @@ struct CreateMaskOpLowering
/// Populate the given list with patterns that convert from ArmSVE to LLVM.
void mlir::populateArmSVELegalizeForLLVMExportPatterns(
- LLVMTypeConverter &converter, RewritePatternSet &patterns) {
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns) {
// Populate conversion patterns
// clang-format off
diff --git a/mlir/lib/Dialect/Bufferization/Transforms/Bufferize.cpp b/mlir/lib/Dialect/Bufferization/Transforms/Bufferize.cpp
index e422086c9fde6d..875d8c40e92cc1 100644
--- a/mlir/lib/Dialect/Bufferization/Transforms/Bufferize.cpp
+++ b/mlir/lib/Dialect/Bufferization/Transforms/Bufferize.cpp
@@ -130,7 +130,7 @@ class BufferizeToMemrefOp
} // namespace
void mlir::bufferization::populateEliminateBufferizeMaterializationsPatterns(
- BufferizeTypeConverter &typeConverter, RewritePatternSet &patterns) {
+ const BufferizeTypeConverter &typeConverter, RewritePatternSet &patterns) {
patterns.add<BufferizeToTensorOp, BufferizeToMemrefOp>(typeConverter,
patterns.getContext());
}
diff --git a/mlir/lib/Dialect/Func/Transforms/DecomposeCallGraphTypes.cpp b/mlir/lib/Dialect/Func/Transforms/DecomposeCallGraphTypes.cpp
index fa030cb18e035d..357f993710a26a 100644
--- a/mlir/lib/Dialect/Func/Transforms/DecomposeCallGraphTypes.cpp
+++ b/mlir/lib/Dialect/Func/Transforms/DecomposeCallGraphTypes.cpp
@@ -37,7 +37,7 @@ template <typename SourceOp>
class DecomposeCallGraphTypesOpConversionPattern
: public OpConversionPattern<SourceOp> {
public:
- DecomposeCallGraphTypesOpConversionPattern(TypeConverter &typeConverter,
+ DecomposeCallGraphTypesOpConversionPattern(const TypeConverter &typeConverter,
MLIRContext *context,
ValueDecomposer &decomposer,
PatternBenefit benefit = 1)
@@ -188,7 +188,7 @@ struct DecomposeCallGraphTypesForCallOp
} // namespace
void mlir::populateDecomposeCallGraphTypesPatterns(
- MLIRContext *context, TypeConverter &typeConverter,
+ MLIRContext *context, const TypeConverter &typeConverter,
ValueDecomposer &decomposer, RewritePatternSet &patterns) {
patterns
.add<DecomposeCallGraphTypesForCallOp, DecomposeCallGraphTypesForFuncArgs,
diff --git a/mlir/lib/Dialect/Func/Transforms/FuncConversions.cpp b/mlir/lib/Dialect/Func/Transforms/FuncConversions.cpp
index d1f3b56dbed738..eb444d665ff260 100644
--- a/mlir/lib/Dialect/Func/Transforms/FuncConversions.cpp
+++ b/mlir/lib/Dialect/Func/Transforms/FuncConversions.cpp
@@ -44,7 +44,7 @@ struct CallOpSignatureConversion : public OpConversionPattern<CallOp> {
} // namespace
void mlir::populateCallOpTypeConversionPattern(RewritePatternSet &patterns,
- TypeConverter &converter) {
+ const TypeConverter &converter) {
patterns.add<CallOpSignatureConversion>(converter, patterns.getContext());
}
@@ -59,7 +59,7 @@ class BranchOpInterfaceTypeConversion
BranchOpInterface>::OpInterfaceConversionPattern;
BranchOpInterfaceTypeConversion(
- TypeConverter &typeConverter, MLIRContext *ctx,
+ const TypeConverter &typeConverter, MLIRContext *ctx,
function_ref<bool(BranchOpInterface, int)> shouldConvertBranchOperand)
: OpInterfaceConversionPattern(typeConverter, ctx, /*benefit=*/1),
shouldConvertBranchOperand(shouldConvertBranchOperand) {}
@@ -115,14 +115,14 @@ class ReturnOpTypeConversion : public OpConversionPattern<ReturnOp> {
} // namespace
void mlir::populateBranchOpInterfaceTypeConversionPattern(
- RewritePatternSet &patterns, TypeConverter &typeConverter,
+ RewritePatternSet &patterns, const TypeConverter &typeConverter,
function_ref<bool(BranchOpInterface, int)> shouldConvertBranchOperand) {
patterns.add<BranchOpInterfaceTypeConversion>(
typeConverter, patterns.getContext(), shouldConvertBranchOperand);
}
bool mlir::isLegalForBranchOpInterfaceTypeConversionPattern(
- Operation *op, TypeConverter &converter) {
+ Operation *op, const TypeConverter &converter) {
// All successor operands of branch like operations must be rewritten.
if (auto branchOp = dyn_cast<BranchOpInterface>(op)) {
for (int p = 0, e = op->getBlock()->getNumSuccessors(); p < e; ++p) {
@@ -137,14 +137,13 @@ bool mlir::isLegalForBranchOpInterfaceTypeConversionPattern(
return false;
}
-void mlir::populateReturnOpTypeConversionPattern(RewritePatternSet &patterns,
- TypeConverter &typeConverter) {
+void mlir::populateReturnOpTypeConversionPattern(
+ RewritePatternSet &patterns, const TypeConverter &typeConverter) {
patterns.add<ReturnOpTypeConversion>(typeConverter, patterns.getContext());
}
-bool mlir::isLegalForReturnOpTypeConversionPattern(Operation *op,
- TypeConverter &converter,
- bool returnOpAlwaysLegal) {
+bool mlir::isLegalForReturnOpTypeConversionPattern(
+ Operation *op, const TypeConverter &converter, bool returnOpAlwaysLegal) {
// If this is a `return` and the user pass wants to convert/transform across
// function boundaries, then `converter` is invoked to check whether the
// `return` op is legal.
diff --git a/mlir/lib/Dialect/Func/Transforms/OneToNFuncConversions.cpp b/mlir/lib/Dialect/Func/Transforms/OneToNFuncConversions.cpp
index 8489396da7a2c4..3b8982257a9c95 100644
--- a/mlir/lib/Dialect/Func/Transforms/OneToNFuncConversions.cpp
+++ b/mlir/lib/Dialect/Func/Transforms/OneToNFuncConversions.cpp
@@ -72,7 +72,7 @@ class ConvertTypesInFuncReturnOp : public OneToNOpConversionPattern<ReturnOp> {
namespace mlir {
-void populateFuncTypeConversionPatterns(TypeConverter &typeConverter,
+void populateFuncTypeConversionPatterns(const TypeConverter &typeConverter,
RewritePatternSet &patterns) {
patterns.add<
// clang-format off
diff --git a/mlir/lib/Dialect/Math/Transforms/ExtendToSupportedTypes.cpp b/mlir/lib/Dialect/Math/Transforms/ExtendToSupportedTypes.cpp
index 1a9eafec9fdd57..a570ed5118ef0b 100644
--- a/mlir/lib/Dialect/Math/Transforms/ExtendToSupportedTypes.cpp
+++ b/mlir/lib/Dialect/Math/Transforms/ExtendToSupportedTypes.cpp
@@ -32,7 +32,7 @@ using namespace mlir;
namespace {
struct ExtendToSupportedTypesRewritePattern final : ConversionPattern {
- ExtendToSupportedTypesRewritePattern(TypeConverter &converter,
+ ExtendToSupportedTypesRewritePattern(const TypeConverter &converter,
MLIRContext *context)
: ConversionPattern(converter, MatchAnyOpTypeTag{}, 1, context) {}
LogicalResult
@@ -114,7 +114,7 @@ LogicalResult ExtendToSupportedTypesRewritePattern::matchAndRewrite(
}
void mlir::math::populateExtendToSupportedTypesPatterns(
- RewritePatternSet &patterns, TypeConverter &typeConverter) {
+ RewritePatternSet &patterns, const TypeConverter &typeConverter) {
patterns.add<ExtendToSupportedTypesRewritePattern>(typeConverter,
patterns.getContext());
}
diff --git a/mlir/lib/Dialect/MemRef/Transforms/EmulateNarrowType.cpp b/mlir/lib/Dialect/MemRef/Transforms/EmulateNarrowType.cpp
index a45b79194a7580..9efea066a03c85 100644
--- a/mlir/lib/Dialect/MemRef/Transforms/EmulateNarrowType.cpp
+++ b/mlir/lib/Dialect/MemRef/Transforms/EmulateNarrowType.cpp
@@ -583,7 +583,7 @@ struct ConvertMemRefExpandShape final
//===----------------------------------------------------------------------===//
void memref::populateMemRefNarrowTypeEmulationPatterns(
- arith::NarrowTypeEmulationConverter &typeConverter,
+ const arith::NarrowTypeEmulationConverter &typeConverter,
RewritePatternSet &patterns) {
// Populate `memref.*` conversion patterns.
diff --git a/mlir/lib/Dialect/MemRef/Transforms/EmulateWideInt.cpp b/mlir/lib/Dialect/MemRef/Transforms/EmulateWideInt.cpp
index 57f0141c95dc56..bc4535f97acf04 100644
--- a/mlir/lib/Dialect/MemRef/Transforms/EmulateWideInt.cpp
+++ b/mlir/lib/Dialect/MemRef/Transforms/EmulateWideInt.cpp
@@ -138,7 +138,7 @@ struct EmulateWideIntPass final
//===----------------------------------------------------------------------===//
void memref::populateMemRefWideIntEmulationPatterns(
- arith::WideIntEmulationConverter &typeConverter,
+ const arith::WideIntEmulationConverter &typeConverter,
RewritePatternSet &patterns) {
// Populate `memref.*` conversion patterns.
patterns.add<ConvertMemRefAlloc, ConvertMemRefLoad, ConvertMemRefStore>(
diff --git a/mlir/lib/Dialect/SCF/Transforms/OneToNTypeConversion.cpp b/mlir/lib/Dialect/SCF/Transforms/OneToNTypeConversion.cpp
index 5aa35e79babfce..4cd17f77dfb941 100644
--- a/mlir/lib/Dialect/SCF/Transforms/OneToNTypeConversion.cpp
+++ b/mlir/lib/Dialect/SCF/Transforms/OneToNTypeConversion.cpp
@@ -198,8 +198,8 @@ class ConvertTypesInSCFForOp final : public OneToNOpConversionPattern<ForOp> {
namespace mlir {
namespace scf {
-void populateSCFStructuralOneToNTypeConversions(TypeConverter &typeConverter,
- RewritePatternSet &patterns) {
+void populateSCFStructuralOneToNTypeConversions(
+ const TypeConverter &typeConverter, RewritePatternSet &patterns) {
patterns.add<
// clang-format off
ConvertTypesInSCFConditionOp,
diff --git a/mlir/lib/Dialect/SCF/Transforms/StructuralTypeConversions.cpp b/mlir/lib/Dialect/SCF/Transforms/StructuralTypeConversions.cpp
index e2cc5b4c5ff49b..93a78056db1944 100644
--- a/mlir/lib/Dialect/SCF/Transforms/StructuralTypeConversions.cpp
+++ b/mlir/lib/Dialect/SCF/Transforms/StructuralTypeConversions.cpp
@@ -248,7 +248,7 @@ class ConvertConditionOpTypes : public OpConversionPattern<ConditionOp> {
} // namespace
void mlir::scf::populateSCFStructuralTypeConversions(
- TypeConverter &typeConverter, RewritePatternSet &patterns) {
+ const TypeConverter &typeConverter, RewritePatternSet &patterns) {
patterns.add<ConvertForOpTypes, ConvertIfOpTypes, ConvertYieldOpTypes,
ConvertWhileOpTypes, ConvertConditionOpTypes>(
typeConverter, patterns.getContext());
@@ -271,7 +271,7 @@ void mlir::scf::populateSCFStructuralTypeConversionTarget(
}
void mlir::scf::populateSCFStructuralTypeConversionsAndLegality(
- TypeConverter &typeConverter, RewritePatternSet &patterns,
+ const TypeConverter &typeConverter, RewritePatternSet &patterns,
ConversionTarget &target) {
populateSCFStructuralTypeConversions(typeConverter, patterns);
populateSCFStructuralTypeConversionTarget(typeConverter, target);
diff --git a/mlir/lib/Dialect/SPIRV/Transforms/SPIRVConversion.cpp b/mlir/lib/Dialect/SPIRV/Transforms/SPIRVConversion.cpp
index d833ec9309baaf..656090314d650e 100644
--- a/mlir/lib/Dialect/SPIRV/Transforms/SPIRVConversion.cpp
+++ b/mlir/lib/Dialect/SPIRV/Transforms/SPIRVConversion.cpp
@@ -1573,8 +1573,8 @@ bool SPIRVConversionTarget::isLegalOp(Operation *op) {
// Public functions for populating patterns
//===----------------------------------------------------------------------===//
-void mlir::populateBuiltinFuncToSPIRVPatterns(SPIRVTypeConverter &typeConverter,
- RewritePatternSet &patterns) {
+void mlir::populateBuiltinFuncToSPIRVPatterns(
+ const SPIRVTypeConverter &typeConverter, RewritePatternSet &patterns) {
patterns.add<FuncOpConversion>(typeConverter, patterns.getContext());
}
diff --git a/mlir/lib/Dialect/SparseTensor/Transforms/SparseIterationToScf.cpp b/mlir/lib/Dialect/SparseTensor/Transforms/SparseIterationToScf.cpp
index 71a229bea990c0..04466d198b5b67 100644
--- a/mlir/lib/Dialect/SparseTensor/Transforms/SparseIterationToScf.cpp
+++ b/mlir/lib/Dialect/SparseTensor/Transforms/SparseIterationToScf.cpp
@@ -434,7 +434,7 @@ mlir::SparseIterationTypeConverter::SparseIterationTypeConverter() {
}
void mlir::populateLowerSparseIterationToSCFPatterns(
- TypeConverter &converter, RewritePatternSet &patterns) {
+ const TypeConverter &converter, RewritePatternSet &patterns) {
IterateOp::getCanonicalizationPatterns(patterns, patterns.getContext());
patterns.add<ExtractIterSpaceConverter, ExtractValOpConverter,
diff --git a/mlir/lib/Dialect/SparseTensor/Transforms/SparseStorageSpecifierToLLVM.cpp b/mlir/lib/Dialect/SparseTensor/Transforms/SparseStorageSpecifierToLLVM.cpp
index 91dad050b28852..d4ddde8afec2ac 100644
--- a/mlir/lib/Dialect/SparseTensor/Transforms/SparseStorageSpecifierToLLVM.cpp
+++ b/mlir/lib/Dialect/SparseTensor/Transforms/SparseStorageSpecifierToLLVM.cpp
@@ -350,8 +350,8 @@ struct StorageSpecifierInitOpConverter
// Public method for populating conversion rules.
//===----------------------------------------------------------------------===//
-void mlir::populateStorageSpecifierToLLVMPatterns(TypeConverter &converter,
- RewritePatternSet &patterns) {
+void mlir::populateStorageSpecifierToLLVMPatterns(
+ const TypeConverter &converter, RewritePatternSet &patterns) {
patterns.add<StorageSpecifierGetOpConverter, StorageSpecifierSetOpConverter,
StorageSpecifierInitOpConverter>(converter,
patterns.getContext());
diff --git a/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorCodegen.cpp b/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorCodegen.cpp
index 164e722c45dba8..062a0ea6cc47cb 100644
--- a/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorCodegen.cpp
+++ b/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorCodegen.cpp
@@ -746,8 +746,8 @@ class SparseTensorAllocConverter
: public OpConversionPattern<bufferization::AllocTensorOp> {
public:
using OpConversionPattern::OpConversionPattern;
- SparseTensorAllocConverter(TypeConverter &typeConverter, MLIRContext *context,
- bool enableInit)
+ SparseTensorAllocConverter(const TypeConverter &typeConverter,
+ MLIRContext *context, bool enableInit)
: OpConversionPattern(typeConverter, context),
enableBufferInitialization(enableInit) {}
@@ -808,8 +808,8 @@ class SparseTensorAllocConverter
class SparseTensorEmptyConverter : public OpConversionPattern<tensor::EmptyOp> {
public:
using OpConversionPattern::OpConversionPattern;
- SparseTensorEmptyConverter(TypeConverter &typeConverter, MLIRContext *context,
- bool enableInit)
+ SparseTensorEmptyConverter(const TypeConverter &typeConverter,
+ MLIRContext *context, bool enableInit)
: OpConversionPattern(typeConverter, context),
enableBufferInitialization(enableInit) {}
@@ -850,7 +850,7 @@ class SparseTensorDeallocConverter
: public OpConversionPattern<bufferization::DeallocTensorOp> {
public:
using OpConversionPattern::OpConversionPattern;
- SparseTensorDeallocConverter(TypeConverter &typeConverter,
+ SparseTensorDeallocConverter(const TypeConverter &typeConverter,
MLIRContext *context, bool createDeallocs)
: OpConversionPattern(typeConverter, context),
createDeallocs(createDeallocs) {}
@@ -1411,7 +1411,7 @@ struct SparseAssembleOpConverter : public OpConversionPattern<AssembleOp> {
struct SparseDisassembleOpConverter
: public OpConversionPattern<DisassembleOp> {
using OpConversionPattern::OpConversionPattern;
- SparseDisassembleOpConverter(TypeConverter &typeConverter,
+ SparseDisassembleOpConverter(const TypeConverter &typeConverter,
MLIRContext *context)
: OpConversionPattern(typeConverter, context) {}
@@ -1604,7 +1604,7 @@ struct SparseHasRuntimeLibraryConverter
/// Populates the given patterns list with conversion rules required for
/// the sparsification of linear algebra operations.
void mlir::populateSparseTensorCodegenPatterns(
- TypeConverter &typeConverter, RewritePatternSet &patterns,
+ const TypeConverter &typeConverter, RewritePatternSet &patterns,
bool createSparseDeallocs, bool enableBufferInitialization) {
patterns.add<
SparseAssembleOpConverter, SparseDisassembleOpConverter,
diff --git a/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorConversion.cpp b/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorConversion.cpp
index f0d162bdb84d96..9ffa64dc821d8f 100644
--- a/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorConversion.cpp
+++ b/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorConversion.cpp
@@ -909,8 +909,8 @@ mlir::SparseTensorTypeToPtrConverter::SparseTensorTypeToPtrConverter() {
/// Populates the given patterns list with conversion rules required for
/// the sparsification of linear algebra operations.
-void mlir::populateSparseTensorConversionPatterns(TypeConverter &typeConverter,
- RewritePatternSet &patterns) {
+void mlir::populateSparseTensorConversionPatterns(
+ const TypeConverter &typeConverter, RewritePatternSet &patterns) {
patterns
.add<SparseReturnConverter, SparseTensorLvlOpConverter,
SparseCastConverter, SparseReMapConverter, SparseTensorNewConverter,
diff --git a/mlir/lib/Dialect/Vector/Transforms/VectorEmulateNarrowType.cpp b/mlir/lib/Dialect/Vector/Transforms/VectorEmulateNarrowType.cpp
index d3296ee38c2496..66362d3ca70fb6 100644
--- a/mlir/lib/Dialect/Vector/Transforms/VectorEmulateNarrowType.cpp
+++ b/mlir/lib/Dialect/Vector/Transforms/VectorEmulateNarrowType.cpp
@@ -1234,7 +1234,7 @@ struct RewriteVectorTranspose : OpRewritePattern<vector::TransposeOp> {
//===----------------------------------------------------------------------===//
void vector::populateVectorNarrowTypeEmulationPatterns(
- arith::NarrowTypeEmulationConverter &typeConverter,
+ const arith::NarrowTypeEmulationConverter &typeConverter,
RewritePatternSet &patterns) {
// Populate `vector.*` conversion patterns.
diff --git a/mlir/lib/Dialect/Vector/Transforms/VectorLinearize.cpp b/mlir/lib/Dialect/Vector/Transforms/VectorLinearize.cpp
index 11917ac1e40226..757631944f224f 100644
--- a/mlir/lib/Dialect/Vector/Transforms/VectorLinearize.cpp
+++ b/mlir/lib/Dialect/Vector/Transforms/VectorLinearize.cpp
@@ -500,7 +500,7 @@ void mlir::vector::populateVectorLinearizeTypeConversionsAndLegality(
}
void mlir::vector::populateVectorLinearizeShuffleLikeOpsPatterns(
- TypeConverter &typeConverter, RewritePatternSet &patterns,
+ const TypeConverter &typeConverter, RewritePatternSet &patterns,
ConversionTarget &target, unsigned int targetBitWidth) {
target.addDynamicallyLegalOp<vector::ShuffleOp>(
[=](vector::ShuffleOp shuffleOp) -> bool {
diff --git a/mlir/lib/Dialect/X86Vector/Transforms/LegalizeForLLVMExport.cpp b/mlir/lib/Dialect/X86Vector/Transforms/LegalizeForLLVMExport.cpp
index c33304c18fe48a..e918473cae9e3a 100644
--- a/mlir/lib/Dialect/X86Vector/Transforms/LegalizeForLLVMExport.cpp
+++ b/mlir/lib/Dialect/X86Vector/Transforms/LegalizeForLLVMExport.cpp
@@ -37,7 +37,7 @@ namespace {
/// results of multi-result intrinsic ops.
template <typename OpTy, typename Intr32OpTy, typename Intr64OpTy>
struct LowerToIntrinsic : public OpConversionPattern<OpTy> {
- explicit LowerToIntrinsic(LLVMTypeConverter &converter)
+ explicit LowerToIntrinsic(const LLVMTypeConverter &converter)
: OpConversionPattern<OpTy>(converter, &converter.getContext()) {}
const LLVMTypeConverter &getTypeConverter() const {
@@ -135,7 +135,7 @@ template <typename... Args>
struct RegistryImpl {
/// Registers the patterns specializing the "main" op to one of the
/// "intrinsic" ops depending on elemental type.
- static void registerPatterns(LLVMTypeConverter &converter,
+ static void registerPatterns(const LLVMTypeConverter &converter,
RewritePatternSet &patterns) {
patterns
.add<LowerToIntrinsic<typename Args::MainOp, typename Args::Intr32Op,
@@ -159,7 +159,7 @@ using Registry = RegistryImpl<
/// Populate the given list with patterns that convert from X86Vector to LLVM.
void mlir::populateX86VectorLegalizeForLLVMExportPatterns(
- LLVMTypeConverter &converter, RewritePatternSet &patterns) {
+ const LLVMTypeConverter &converter, RewritePatternSet &patterns) {
Registry::registerPatterns(converter, patterns);
patterns.add<MaskCompressOpConversion, RsqrtOpConversion, DotOpConversion>(
converter);
diff --git a/mlir/lib/Transforms/Utils/OneToNTypeConversion.cpp b/mlir/lib/Transforms/Utils/OneToNTypeConversion.cpp
index f6e8e9e7ad339e..19e29d48623e04 100644
--- a/mlir/lib/Transforms/Utils/OneToNTypeConversion.cpp
+++ b/mlir/lib/Transforms/Utils/OneToNTypeConversion.cpp
@@ -418,7 +418,7 @@ class FunctionOpInterfaceSignatureConversion : public OneToNConversionPattern {
public:
FunctionOpInterfaceSignatureConversion(StringRef functionLikeOpName,
MLIRContext *ctx,
- TypeConverter &converter)
+ const TypeConverter &converter)
: OneToNConversionPattern(converter, functionLikeOpName, /*benefit=*/1,
ctx) {}
@@ -466,7 +466,7 @@ class FunctionOpInterfaceSignatureConversion : public OneToNConversionPattern {
} // namespace
void populateOneToNFunctionOpInterfaceTypeConversionPattern(
- StringRef functionLikeOpName, TypeConverter &converter,
+ StringRef functionLikeOpName, const TypeConverter &converter,
RewritePatternSet &patterns) {
patterns.add<FunctionOpInterfaceSignatureConversion>(
functionLikeOpName, patterns.getContext(), converter);
diff --git a/mlir/test/lib/Conversion/OneToNTypeConversion/TestOneToNTypeConversionPass.cpp b/mlir/test/lib/Conversion/OneToNTypeConversion/TestOneToNTypeConversionPass.cpp
index cc1af59c5e15bb..1ea65109bf79db 100644
--- a/mlir/test/lib/Conversion/OneToNTypeConversion/TestOneToNTypeConversionPass.cpp
+++ b/mlir/test/lib/Conversion/OneToNTypeConversion/TestOneToNTypeConversionPass.cpp
@@ -129,8 +129,9 @@ class ConvertGetTupleElementOp
} // namespace
-static void populateDecomposeTuplesTestPatterns(TypeConverter &typeConverter,
- RewritePatternSet &patterns) {
+static void
+populateDecomposeTuplesTestPatterns(const TypeConverter &typeConverter,
+ RewritePatternSet &patterns) {
patterns.add<
// clang-format off
ConvertMakeTupleOp,
More information about the flang-commits
mailing list