[mlir][Conversion] Store const type converter in ConversionPattern

[springerm]

Summary

ConversionPatterns do not (and should not) modify the type converter that they are using.

• Make ConversionPattern::typeConverter const.
• Make member functions of the LLVMTypeConverter const.
• Conversion patterns take a const type converter.
• Various helper functions (that are called from patterns) now also take a const type converter.

Depends On: D157611

[nicolasvasilache]

Great, thanks for untangling this, I had a similar thing in flight but did not get all the spots yet so happy to use this!

[springerm]

fix flang

[kuhar]

What is the motivation for this change beyond cleaning up the code? IIRC the mlir policy discourages const in the general case: https://mlir.llvm.org/getting_started/DeveloperGuide/#style-guide, so I'm wondering if you plan some future API changes stacked on top of this one.

Do we need this because matchAndRewrite is marked as const? As an alternative could we remove the const qualifiers in RewritePattern, or has the ship already sailed there?

[zero9178]

In D157601#4577758, @kuhar wrote:

What is the motivation for this change beyond cleaning up the code? IIRC the mlir policy discourages const in the general case: https://mlir.llvm.org/getting_started/DeveloperGuide/#style-guide, so I'm wondering if you plan some future API changes stacked on top of this one.

Do we need this because matchAndRewrite is marked as const? As an alternative could we remove the const qualifiers in RewritePattern, or has the ship already sailed there?

The style guide reads to me like this only applies to IR constructs. The linked rationale doc even clarifies further and seems to suggest that it only applies to operations, blocks and regions. This has previously come up here: https://discourse.llvm.org/t/clarifying-usage-of-const/4986/2. So I don't think this patch is in conflict with that policy.

[kuhar]

In D157601#4577815, @zero9178 wrote:

In D157601#4577758, @kuhar wrote:

What is the motivation for this change beyond cleaning up the code? IIRC the mlir policy discourages const in the general case: https://mlir.llvm.org/getting_started/DeveloperGuide/#style-guide, so I'm wondering if you plan some future API changes stacked on top of this one.

Do we need this because matchAndRewrite is marked as const? As an alternative could we remove the const qualifiers in RewritePattern, or has the ship already sailed there?

The style guide reads to me like this only applies to IR constructs. The linked rationale doc even clarifies further and seems to suggest that it only applies to operations, blocks and regions. This has previously come up here: https://discourse.llvm.org/t/clarifying-usage-of-const/4986/2. So I don't think this patch is in conflict with that policy.

Ah, I vaguely recall @Mogball arguing the opposite interpretation when it came to local variables etc. If that's the case, I guess we fallback to general rules of idiomatic C++ in the rest of the project. Separately from this revision, we may expand on the constness policy on the website because this keeps coming up.

[springerm]

In D157601#4577758, @kuhar wrote:

What is the motivation for this change beyond cleaning up the code?

Yes, this is a cleanup. And it hardens the API to prevent users from changing the type converter while a dialect conversion is in progress. It is not related to matchAndRewrite (or the fact that it is const).
Discussion: https://discourse.llvm.org/t/rfc-almost-all-uses-of-typeconverter-should-be-const/72689

[kuhar]

In D157601#4577826, @springerm wrote:

In D157601#4577758, @kuhar wrote:

What is the motivation for this change beyond cleaning up the code?

Yes, this is a cleanup. And it hardens the API to prevent users from changing the type converter while a dialect conversion is in progress. It is not related to matchAndRewrite (or the fact that it is const).
Discussion: https://discourse.llvm.org/t/rfc-almost-all-uses-of-typeconverter-should-be-const/72689

SGTM. Could you add the link to the RFC to the commit description?

[Mogball]

Apologies for missing the RFC, but I am a small -1 on this direction. The rationale I have for allowing mutability of a type converter instance is caching expensive type conversions in a hashmap inside the type converter. It's specific to my use case and I can always add a separate path, however.

On the general const policy of the codebase, the emphasis is definitely on the core language constructs but the overall coding style of MLIR is not go crazy and mark things as const unless there is an invariant to enforce. It seems with TypeConverter the invariant is useful to enforce so this makes sense to me.

[springerm]

In D157601#4577837, @Mogball wrote:

The rationale I have for allowing mutability of a type converter instance is caching expensive type conversions in a hashmap inside the type converter. It's specific to my use case and I can always add a separate path, however.

I made sure that we can still cache type conversions: D157597

[springerm]

rebase

[mehdi_amini]

In D157601#4577843, @springerm wrote:

In D157601#4577837, @Mogball wrote:

The rationale I have for allowing mutability of a type converter instance is caching expensive type conversions in a hashmap inside the type converter. It's specific to my use case and I can always add a separate path, however.

I made sure that we can still cache type conversions: D157597

Seems like everyone is happy with this?
If so, thanks for the cleanup: mutable state is harder to reason through!

mlir/include/mlir/Conversion/LLVMCommon/MemRefBuilder.h

public:
   /// Construct a helper for the given descriptor value.
   explicit MemRefDescriptor(Value descriptor);
   /// Builds IR creating an `undef` value of the descriptor type.
   static MemRefDescriptor undef(OpBuilder &builder, Location loc,
                                 Type descriptorType);
   /// Builds IR creating a MemRef descriptor that represents `type` and
   /// populates it with static shape and stride information extracted from the
   /// type.
-  static MemRefDescriptor fromStaticShape(OpBuilder &builder, Location loc,
-                                          LLVMTypeConverter &typeConverter,
-                                          MemRefType type, Value memory);
-  static MemRefDescriptor fromStaticShape(OpBuilder &builder, Location loc,
-                                          LLVMTypeConverter &typeConverter,
-                                          MemRefType type, Value memory,
-                                          Value alignedMemory);
+  static MemRefDescriptor
+  fromStaticShape(OpBuilder &builder, Location loc,
+                  const LLVMTypeConverter &typeConverter, MemRefType type,
+                  Value memory);
+  static MemRefDescriptor
+  fromStaticShape(OpBuilder &builder, Location loc,
+                  const LLVMTypeConverter &typeConverter, MemRefType type,
+                  Value memory, Value alignedMemory);
 
   /// Builds IR extracting the allocated pointer from the descriptor.
   Value allocatedPtr(OpBuilder &builder, Location loc);
   /// Builds IR inserting the allocated pointer into the descriptor.
   void setAllocatedPtr(OpBuilder &builder, Location loc, Value ptr);
 
   /// Builds IR extracting the aligned pointer from the descriptor.
   Value alignedPtr(OpBuilder &builder, Location loc);
public:
   LLVM::LLVMPointerType getElementPtrType();
 
   /// Builds IR for getting the start address of the buffer represented
   /// by this memref:
   /// `memref.alignedPtr + memref.offset * sizeof(type.getElementType())`.
   /// \note there is no setter for this one since it is derived from alignedPtr
   /// and offset.
   Value bufferPtr(OpBuilder &builder, Location loc,
-                  LLVMTypeConverter &converter, MemRefType type);
+                  const LLVMTypeConverter &converter, MemRefType type);
 
   /// Builds IR populating a MemRef descriptor structure from a list of
   /// individual values composing that descriptor, in the following order:
   /// - allocated pointer;
   /// - aligned pointer;
   /// - offset;
   /// - <rank> sizes;
   /// - <rank> shapes;
   /// where <rank> is the MemRef rank as provided in `type`.
   static Value pack(OpBuilder &builder, Location loc,
-                    LLVMTypeConverter &converter, MemRefType type,
+                    const LLVMTypeConverter &converter, MemRefType type,
                     ValueRange values);
 
   /// Builds IR extracting individual elements of a MemRef descriptor structure
   /// and returning them as `results` list.
   static void unpack(OpBuilder &builder, Location loc, Value packed,
                      MemRefType type, SmallVectorImpl<Value> &results);
 
   /// Returns the number of non-aggregate values that would be produced by
public:
   /// Builds IR setting ranked memref descriptor ptr
   void setMemRefDescPtr(OpBuilder &builder, Location loc, Value value);
 
   /// Builds IR populating an unranked MemRef descriptor structure from a list
   /// of individual constituent values in the following order:
   /// - rank of the memref;
   /// - pointer to the memref descriptor.
   static Value pack(OpBuilder &builder, Location loc,
-                    LLVMTypeConverter &converter, UnrankedMemRefType type,
+                    const LLVMTypeConverter &converter, UnrankedMemRefType type,
                     ValueRange values);
 
   /// Builds IR extracting individual elements that compose an unranked memref
   /// descriptor and returns them as `results` list.
   static void unpack(OpBuilder &builder, Location loc, Value packed,
                      SmallVectorImpl<Value> &results);
 
   /// Returns the number of non-aggregate values that would be produced by
   /// `unpack`.
   static unsigned getNumUnpackedValues() { return 2; }
 
   /// Builds IR computing the sizes in bytes (suitable for opaque allocation)
   /// and appends the corresponding values into `sizes`. `addressSpaces`
   /// which must have the same length as `values`, is needed to handle layouts
   /// where sizeof(ptr addrspace(N)) != sizeof(ptr addrspace(0)).
   static void computeSizes(OpBuilder &builder, Location loc,
-                           LLVMTypeConverter &typeConverter,
+                           const LLVMTypeConverter &typeConverter,
                            ArrayRef<UnrankedMemRefDescriptor> values,
                            ArrayRef<unsigned> addressSpaces,
                            SmallVectorImpl<Value> &sizes);
 
   /// TODO: The following accessors don't take alignment rules between elements
   /// of the descriptor struct into account. For some architectures, it might be
   /// necessary to extend them and to use `llvm::DataLayout` contained in
   /// `LLVMTypeConverter`.
 
   /// Builds IR extracting the allocated pointer from the descriptor.
   static Value allocatedPtr(OpBuilder &builder, Location loc,
                             Value memRefDescPtr,
                             LLVM::LLVMPointerType elemPtrType);
   /// Builds IR inserting the allocated pointer into the descriptor.
   static void setAllocatedPtr(OpBuilder &builder, Location loc,
                               Value memRefDescPtr,
                               LLVM::LLVMPointerType elemPtrType,
                               Value allocatedPtr);
 
   /// Builds IR extracting the aligned pointer from the descriptor.
   static Value alignedPtr(OpBuilder &builder, Location loc,
-                          LLVMTypeConverter &typeConverter, Value memRefDescPtr,
+                          const LLVMTypeConverter &typeConverter,
+                          Value memRefDescPtr,
                           LLVM::LLVMPointerType elemPtrType);
   /// Builds IR inserting the aligned pointer into the descriptor.
   static void setAlignedPtr(OpBuilder &builder, Location loc,
-                            LLVMTypeConverter &typeConverter,
+                            const LLVMTypeConverter &typeConverter,
                             Value memRefDescPtr,
                             LLVM::LLVMPointerType elemPtrType,
                             Value alignedPtr);
 
   /// Builds IR for getting the pointer to the offset's location.
   /// Returns a pointer to a convertType(index), which points to the beggining
   /// of a struct {index, index[rank], index[rank]}.
   static Value offsetBasePtr(OpBuilder &builder, Location loc,
-                             LLVMTypeConverter &typeConverter,
+                             const LLVMTypeConverter &typeConverter,
                              Value memRefDescPtr,
                              LLVM::LLVMPointerType elemPtrType);
   /// Builds IR extracting the offset from the descriptor.
   static Value offset(OpBuilder &builder, Location loc,
-                      LLVMTypeConverter &typeConverter, Value memRefDescPtr,
-                      LLVM::LLVMPointerType elemPtrType);
+                      const LLVMTypeConverter &typeConverter,
+                      Value memRefDescPtr, LLVM::LLVMPointerType elemPtrType);
   /// Builds IR inserting the offset into the descriptor.
   static void setOffset(OpBuilder &builder, Location loc,
-                        LLVMTypeConverter &typeConverter, Value memRefDescPtr,
-                        LLVM::LLVMPointerType elemPtrType, Value offset);
+                        const LLVMTypeConverter &typeConverter,
+                        Value memRefDescPtr, LLVM::LLVMPointerType elemPtrType,
+                        Value offset);
 
   /// Builds IR extracting the pointer to the first element of the size array.
   static Value sizeBasePtr(OpBuilder &builder, Location loc,
-                           LLVMTypeConverter &typeConverter,
+                           const LLVMTypeConverter &typeConverter,
                            Value memRefDescPtr,
                            LLVM::LLVMPointerType elemPtrType);
   /// Builds IR extracting the size[index] from the descriptor.
   static Value size(OpBuilder &builder, Location loc,
-                    LLVMTypeConverter &typeConverter, Value sizeBasePtr,
+                    const LLVMTypeConverter &typeConverter, Value sizeBasePtr,
                     Value index);
   /// Builds IR inserting the size[index] into the descriptor.
   static void setSize(OpBuilder &builder, Location loc,
-                      LLVMTypeConverter &typeConverter, Value sizeBasePtr,
+                      const LLVMTypeConverter &typeConverter, Value sizeBasePtr,
                       Value index, Value size);
 
   /// Builds IR extracting the pointer to the first element of the stride array.
   static Value strideBasePtr(OpBuilder &builder, Location loc,
-                             LLVMTypeConverter &typeConverter,
+                             const LLVMTypeConverter &typeConverter,
                              Value sizeBasePtr, Value rank);
   /// Builds IR extracting the stride[index] from the descriptor.
   static Value stride(OpBuilder &builder, Location loc,
-                      LLVMTypeConverter &typeConverter, Value strideBasePtr,
-                      Value index, Value stride);
+                      const LLVMTypeConverter &typeConverter,
+                      Value strideBasePtr, Value index, Value stride);
   /// Builds IR inserting the stride[index] into the descriptor.
   static void setStride(OpBuilder &builder, Location loc,
-                        LLVMTypeConverter &typeConverter, Value strideBasePtr,
-                        Value index, Value stride);
+                        const LLVMTypeConverter &typeConverter,
+                        Value strideBasePtr, Value index, Value stride);
 };
 
 } // namespace mlir
 
 #endif // MLIR_CONVERSION_LLVMCOMMON_MEMREFBUILDER_H

mlir/include/mlir/Conversion/LLVMCommon/Pattern.h

 
 namespace LLVM {
 namespace detail {
 /// Replaces the given operation "op" with a new operation of type "targetOp"
 /// and given operands.
 LogicalResult oneToOneRewrite(Operation *op, StringRef targetOp,
                               ValueRange operands,
                               ArrayRef<NamedAttribute> targetAttrs,
-                              LLVMTypeConverter &typeConverter,
+                              const LLVMTypeConverter &typeConverter,
                               ConversionPatternRewriter &rewriter);
 
 } // namespace detail
 } // namespace LLVM
 
 /// Base class for operation conversions targeting the LLVM IR dialect. It
 /// provides the conversion patterns with access to the LLVMTypeConverter and
 /// the LowerToLLVMOptions. The class captures the LLVMTypeConverter and the
 /// LowerToLLVMOptions by reference meaning the references have to remain alive
 /// during the entire pattern lifetime.
 class ConvertToLLVMPattern : public ConversionPattern {
 public:
   ConvertToLLVMPattern(StringRef rootOpName, MLIRContext *context,
-                       LLVMTypeConverter &typeConverter,
+                       const LLVMTypeConverter &typeConverter,
                        PatternBenefit benefit = 1);
 
 protected:
   /// Returns the LLVM dialect.
   LLVM::LLVMDialect &getDialect() const;
 
-  LLVMTypeConverter *getTypeConverter() const;
+  const LLVMTypeConverter *getTypeConverter() const;
 
   /// Gets the MLIR type wrapping the LLVM integer type whose bit width is
   /// defined by the used type converter.
   Type getIndexType() const;
 
   /// Gets the MLIR type wrapping the LLVM integer type whose bit width
   /// corresponds to that of a LLVM pointer type.
   Type getIntPtrType(unsigned addressSpace = 0) const;

 
 /// Utility class for operation conversions targeting the LLVM dialect that
 /// match exactly one source operation.
 template <typename SourceOp>
 class ConvertOpToLLVMPattern : public ConvertToLLVMPattern {
 public:
   using OpAdaptor = typename SourceOp::Adaptor;
 
-  explicit ConvertOpToLLVMPattern(LLVMTypeConverter &typeConverter,
+  explicit ConvertOpToLLVMPattern(const LLVMTypeConverter &typeConverter,
                                   PatternBenefit benefit = 1)
       : ConvertToLLVMPattern(SourceOp::getOperationName(),
                              &typeConverter.getContext(), typeConverter,
                              benefit) {}
 
   /// Wrappers around the RewritePattern methods that pass the derived op type.
   void rewrite(Operation *op, ArrayRef<Value> operands,
                ConversionPatternRewriter &rewriter) const final {

mlir/include/mlir/Conversion/LLVMCommon/TypeConverter.h

 class LLVMFunctionType;
 class LLVMStructType;
 } // namespace LLVM
 
 /// Conversion from types to the LLVM IR dialect.
 class LLVMTypeConverter : public TypeConverter {
   /// Give structFuncArgTypeConverter access to memref-specific functions.
   friend LogicalResult
-  structFuncArgTypeConverter(LLVMTypeConverter &converter, Type type,
+  structFuncArgTypeConverter(const LLVMTypeConverter &converter, Type type,
                              SmallVectorImpl<Type> &result);
 
 public:
   using TypeConverter::convertType;
 
   /// Create an LLVMTypeConverter using the default LowerToLLVMOptions.
   /// Optionally takes a data layout analysis to use in conversions.
   LLVMTypeConverter(MLIRContext *ctx,
                     const DataLayoutAnalysis *analysis = nullptr);
 
   /// Create an LLVMTypeConverter using custom LowerToLLVMOptions. Optionally
   /// takes a data layout analysis to use in conversions.
   LLVMTypeConverter(MLIRContext *ctx, const LowerToLLVMOptions &options,
                     const DataLayoutAnalysis *analysis = nullptr);
 
   /// Convert a function type.  The arguments and results are converted one by
   /// one and results are packed into a wrapped LLVM IR structure type. `result`
   /// is populated with argument mapping.
   Type convertFunctionSignature(FunctionType funcTy, bool isVariadic,
                                 bool useBarePtrCallConv,
-                                SignatureConversion &result);
+                                SignatureConversion &result) const;
 
   /// Convert a non-empty list of types to be returned from a function into an
   /// LLVM-compatible type. In particular, if more than one value is returned,
   /// create an LLVM dialect structure type with elements that correspond to
   /// each of the types converted with `convertCallingConventionType`.
   Type packFunctionResults(TypeRange types,
-                           bool useBarePointerCallConv = false);
+                           bool useBarePointerCallConv = false) const;
 
   /// Convert a non-empty list of types of values produced by an operation into
   /// an LLVM-compatible type. In particular, if more than one value is
   /// produced, create a literal structure with elements that correspond to each
   /// of the LLVM-compatible types converted with `convertType`.
-  Type packOperationResults(TypeRange types);
+  Type packOperationResults(TypeRange types) const;
 
   /// Convert a type in the context of the default or bare pointer calling
   /// convention. Calling convention sensitive types, such as MemRefType and
   /// UnrankedMemRefType, are converted following the specific rules for the
   /// calling convention. Calling convention independent types are converted
   /// following the default LLVM type conversions.
   Type convertCallingConventionType(Type type,
-                                    bool useBarePointerCallConv = false);
+                                    bool useBarePointerCallConv = false) const;
 
   /// Promote the bare pointers in 'values' that resulted from memrefs to
   /// descriptors. 'stdTypes' holds the types of 'values' before the conversion
   /// to the LLVM-IR dialect (i.e., MemRefType, or any other builtin type).
   void promoteBarePtrsToDescriptors(ConversionPatternRewriter &rewriter,
                                     Location loc, ArrayRef<Type> stdTypes,
-                                    SmallVectorImpl<Value> &values);
+                                    SmallVectorImpl<Value> &values) const;
 
   /// Returns the MLIR context.
-  MLIRContext &getContext();
+  MLIRContext &getContext() const;
 
   /// Returns the LLVM dialect.
-  LLVM::LLVMDialect *getDialect() { return llvmDialect; }
+  LLVM::LLVMDialect *getDialect() const { return llvmDialect; }
 
   const LowerToLLVMOptions &getOptions() const { return options; }
 
   /// Set the lowering options to `newOptions`. Note: using this after some
   /// some conversions have been performed can lead to inconsistencies in the
   /// IR.
   void dangerousSetOptions(LowerToLLVMOptions newOptions) {
     options = std::move(newOptions);
   }
 
   /// Promote the LLVM representation of all operands including promoting MemRef
   /// descriptors to stack and use pointers to struct to avoid the complexity
   /// of the platform-specific C/C++ ABI lowering related to struct argument
   /// passing.
   SmallVector<Value, 4> promoteOperands(Location loc, ValueRange opOperands,
                                         ValueRange operands, OpBuilder &builder,
-                                        bool useBarePtrCallConv = false);
+                                        bool useBarePtrCallConv = false) const;
 
   /// Promote the LLVM struct representation of one MemRef descriptor to stack
   /// and use pointer to struct to avoid the complexity of the platform-specific
   /// C/C++ ABI lowering related to struct argument passing.
   Value promoteOneMemRefDescriptor(Location loc, Value operand,
-                                   OpBuilder &builder);
+                                   OpBuilder &builder) const;
 
   /// Converts the function type to a C-compatible format, in particular using
   /// pointers to memref descriptors for arguments. Also converts the return
   /// type to a pointer argument if it is a struct. Returns true if this
   /// was the case.
   std::pair<LLVM::LLVMFunctionType, LLVM::LLVMStructType>
-  convertFunctionTypeCWrapper(FunctionType type);
+  convertFunctionTypeCWrapper(FunctionType type) const;
 
   /// Returns the data layout to use during and after conversion.
-  const llvm::DataLayout &getDataLayout() { return options.dataLayout; }
+  const llvm::DataLayout &getDataLayout() const { return options.dataLayout; }
 
   /// Returns the data layout analysis to query during conversion.
   const DataLayoutAnalysis *getDataLayoutAnalysis() const {
     return dataLayoutAnalysis;
   }
 
   /// Gets the LLVM representation of the index type. The returned type is an
   /// integer type with the size configured for this type converter.
-  Type getIndexType();
+  Type getIndexType() const;
 
   /// Returns true if using opaque pointers was enabled in the lowering options.
   bool useOpaquePointers() const { return getOptions().useOpaquePointers; }
 
   /// Creates an LLVM pointer type with the given element type and address
   /// space.
   /// This function is meant to be used in code supporting both typed and opaque
   /// pointers, as it will create an opaque pointer with the given address space
   /// if opaque pointers are enabled in the lowering options.
   LLVM::LLVMPointerType getPointerType(Type elementType,
-                                       unsigned addressSpace = 0);
+                                       unsigned addressSpace = 0) const;
 
   /// Gets the bitwidth of the index type when converted to LLVM.
-  unsigned getIndexTypeBitwidth() { return options.getIndexBitwidth(); }
+  unsigned getIndexTypeBitwidth() const { return options.getIndexBitwidth(); }
 
   /// Gets the pointer bitwidth.
-  unsigned getPointerBitwidth(unsigned addressSpace = 0);
+  unsigned getPointerBitwidth(unsigned addressSpace = 0) const;
 
   /// Returns the size of the memref descriptor object in bytes.
-  unsigned getMemRefDescriptorSize(MemRefType type, const DataLayout &layout);
+  unsigned getMemRefDescriptorSize(MemRefType type,
+                                   const DataLayout &layout) const;
 
   /// Returns the size of the unranked memref descriptor object in bytes.
   unsigned getUnrankedMemRefDescriptorSize(UnrankedMemRefType type,
-                                           const DataLayout &layout);
+                                           const DataLayout &layout) const;
 
   /// Return the LLVM address space corresponding to the memory space of the
   /// memref type `type` or failure if the memory space cannot be converted to
   /// an integer.
-  FailureOr<unsigned> getMemRefAddressSpace(BaseMemRefType type);
+  FailureOr<unsigned> getMemRefAddressSpace(BaseMemRefType type) const;
 
   /// Check if a memref type can be converted to a bare pointer.
   static bool canConvertToBarePtr(BaseMemRefType type);
 
 protected:
   /// Pointer to the LLVM dialect.
   LLVM::LLVMDialect *llvmDialect;
 
 private:
   /// Convert a function type.  The arguments and results are converted one by
   /// one.  Additionally, if the function returns more than one value, pack the
   /// results into an LLVM IR structure type so that the converted function type
   /// returns at most one result.
-  Type convertFunctionType(FunctionType type);
+  Type convertFunctionType(FunctionType type) const;
 
   /// Convert the index type.  Uses llvmModule data layout to create an integer
   /// of the pointer bitwidth.
-  Type convertIndexType(IndexType type);
+  Type convertIndexType(IndexType type) const;
 
   /// Convert an integer type `i*` to `!llvm<"i*">`.
-  Type convertIntegerType(IntegerType type);
+  Type convertIntegerType(IntegerType type) const;
 
   /// Convert a floating point type: `f16` to `f16`, `f32` to
   /// `f32` and `f64` to `f64`.  `bf16` is not supported
   /// by LLVM. 8-bit float types are converted to 8-bit integers as this is how
   /// all LLVM backends that support them currently represent them.
-  Type convertFloatType(FloatType type);
+  Type convertFloatType(FloatType type) const;
 
   /// Convert complex number type: `complex<f16>` to `!llvm<"{ half, half }">`,
   /// `complex<f32>` to `!llvm<"{ float, float }">`, and `complex<f64>` to
   /// `!llvm<"{ double, double }">`. `complex<bf16>` is not supported.
-  Type convertComplexType(ComplexType type);
+  Type convertComplexType(ComplexType type) const;
 
   /// Convert a memref type into an LLVM type that captures the relevant data.
-  Type convertMemRefType(MemRefType type);
+  Type convertMemRefType(MemRefType type) const;
 
   /// Convert a memref type into a list of LLVM IR types that will form the
   /// memref descriptor. If `unpackAggregates` is true the `sizes` and `strides`
   /// arrays in the descriptors are unpacked to individual index-typed elements,
   /// else they are are kept as rank-sized arrays of index type. In particular,
   /// the list will contain:
   /// - two pointers to the memref element type, followed by
   /// - an index-typed offset, followed by
   /// - (if unpackAggregates = true)
   ///    - one index-typed size per dimension of the memref, followed by
   ///    - one index-typed stride per dimension of the memref.
   /// - (if unpackArrregates = false)
   ///   - one rank-sized array of index-type for the size of each dimension
   ///   - one rank-sized array of index-type for the stride of each dimension
   ///
   /// For example, memref<?x?xf32> is converted to the following list:
   /// - `!llvm<"float*">` (allocated pointer),
   /// - `!llvm<"float*">` (aligned pointer),
   /// - `i64` (offset),
   /// - `i64`, `i64` (sizes),
   /// - `i64`, `i64` (strides).
   /// These types can be recomposed to a memref descriptor struct.
   SmallVector<Type, 5> getMemRefDescriptorFields(MemRefType type,
-                                                 bool unpackAggregates);
+                                                 bool unpackAggregates) const;
 
   /// Convert an unranked memref type into a list of non-aggregate LLVM IR types
   /// that will form the unranked memref descriptor. In particular, this list
   /// contains:
   /// - an integer rank, followed by
   /// - a pointer to the memref descriptor struct.
   /// For example, memref<*xf32> is converted to the following list:
   /// i64 (rank)
   /// !llvm<"i8*"> (type-erased pointer).
   /// These types can be recomposed to a unranked memref descriptor struct.
-  SmallVector<Type, 2> getUnrankedMemRefDescriptorFields();
+  SmallVector<Type, 2> getUnrankedMemRefDescriptorFields() const;
 
   /// Convert an unranked memref type to an LLVM type that captures the
   /// runtime rank and a pointer to the static ranked memref desc
-  Type convertUnrankedMemRefType(UnrankedMemRefType type);
+  Type convertUnrankedMemRefType(UnrankedMemRefType type) const;
 
   /// Convert a memref type to a bare pointer to the memref element type.
-  Type convertMemRefToBarePtr(BaseMemRefType type);
+  Type convertMemRefToBarePtr(BaseMemRefType type) const;
 
   /// Convert a 1D vector type into an LLVM vector type.
-  Type convertVectorType(VectorType type);
+  Type convertVectorType(VectorType type) const;
 
   /// Options for customizing the llvm lowering.
   LowerToLLVMOptions options;
 
   /// Data layout analysis mapping scopes to layouts active in them.
   const DataLayoutAnalysis *dataLayoutAnalysis;
 };
 
 /// Callback to convert function argument types. It converts a MemRef function
 /// argument to a list of non-aggregate types containing descriptor
 /// information, and an UnrankedmemRef function argument to a list containing
 /// the rank and a pointer to a descriptor struct.
-LogicalResult structFuncArgTypeConverter(LLVMTypeConverter &converter,
+LogicalResult structFuncArgTypeConverter(const LLVMTypeConverter &converter,
                                          Type type,
                                          SmallVectorImpl<Type> &result);
 
 /// Callback to convert function argument types. It converts MemRef function
 /// arguments to bare pointers to the MemRef element type.
-LogicalResult barePtrFuncArgTypeConverter(LLVMTypeConverter &converter,
+LogicalResult barePtrFuncArgTypeConverter(const LLVMTypeConverter &converter,
                                           Type type,
                                           SmallVectorImpl<Type> &result);
 
 } // namespace mlir
 
 #endif // MLIR_CONVERSION_LLVMCOMMON_TYPECONVERTER_H

mlir/include/mlir/Conversion/LLVMCommon/VectorPattern.h

struct NDVectorTypeInfo {
   SmallVector<int64_t, 4> arraySizes;
 };
 
 // For >1-D vector types, extracts the necessary information to iterate over all
 // 1-D subvectors in the underlying llrepresentation of the n-D vector
 // Iterates on the llvm array type until we hit a non-array type (which is
 // asserted to be an llvm vector type).
 NDVectorTypeInfo extractNDVectorTypeInfo(VectorType vectorType,
-                                         LLVMTypeConverter &converter);
+                                         const LLVMTypeConverter &converter);
 
 // Express `linearIndex` in terms of coordinates of `basis`.
 // Returns the empty vector when linearIndex is out of the range [0, P] where
 // P is the product of all the basis coordinates.
 //
 // Prerequisites:
 //   Basis is an array of nonnegative integers (signed type inherited from
 //   vector shape type).
 SmallVector<int64_t, 4> getCoordinates(ArrayRef<int64_t> basis,
                                        unsigned linearIndex);
 
 // Iterate of linear index, convert to coords space and insert splatted 1-D
 // vector in each position.
 void nDVectorIterate(const NDVectorTypeInfo &info, OpBuilder &builder,
                      function_ref<void(ArrayRef<int64_t>)> fun);
 
 LogicalResult handleMultidimensionalVectors(
-    Operation *op, ValueRange operands, LLVMTypeConverter &typeConverter,
+    Operation *op, ValueRange operands, const LLVMTypeConverter &typeConverter,
     std::function<Value(Type, ValueRange)> createOperand,
     ConversionPatternRewriter &rewriter);
 
 LogicalResult vectorOneToOneRewrite(Operation *op, StringRef targetOp,
                                     ValueRange operands,
                                     ArrayRef<NamedAttribute> targetAttrs,
-                                    LLVMTypeConverter &typeConverter,
+                                    const LLVMTypeConverter &typeConverter,
                                     ConversionPatternRewriter &rewriter);
 } // namespace detail
 } // namespace LLVM
 
 // Default attribute conversion class, which passes all source attributes
 // through to the target op, unmodified.
 template <typename SourceOp, typename TargetOp>
 class AttrConvertPassThrough {

mlir/include/mlir/Conversion/MemRefToLLVM/AllocLikeConversion.h

 
 /// Lowering for memory allocation ops.
 struct AllocationOpLLVMLowering : public ConvertToLLVMPattern {
   using ConvertToLLVMPattern::createIndexAttrConstant;
   using ConvertToLLVMPattern::getIndexType;
   using ConvertToLLVMPattern::getVoidPtrType;
 
   explicit AllocationOpLLVMLowering(StringRef opName,
-                                    LLVMTypeConverter &converter,
+                                    const LLVMTypeConverter &converter,
                                     PatternBenefit benefit = 1)
       : ConvertToLLVMPattern(opName, &converter.getContext(), converter,
                              benefit) {}
 
 protected:
   /// Computes the aligned value for 'input' as follows:
   ///   bumped = input + alignement - 1
   ///   aligned = bumped - bumped % alignment
private:
 
   /// The minimum alignment to use with aligned_alloc (has to be a power of 2).
   static constexpr uint64_t kMinAlignedAllocAlignment = 16UL;
 };
 
 /// Lowering for AllocOp and AllocaOp.
 struct AllocLikeOpLLVMLowering : public AllocationOpLLVMLowering {
   explicit AllocLikeOpLLVMLowering(StringRef opName,
-                                   LLVMTypeConverter &converter,
+                                   const LLVMTypeConverter &converter,
                                    PatternBenefit benefit = 1)
       : AllocationOpLLVMLowering(opName, converter, benefit) {}
 
 protected:
   /// Allocates the underlying buffer. Returns the allocated pointer and the
   /// aligned pointer.
   virtual std::tuple<Value, Value>
   allocateBuffer(ConversionPatternRewriter &rewriter, Location loc, Value size,

mlir/include/mlir/Dialect/SPIRV/Transforms/SPIRVConversion.h

public:
   Type getIndexType() const;
 
   const spirv::TargetEnv &getTargetEnv() const { return targetEnv; }
 
   /// Returns the options controlling the SPIR-V type converter.
   const SPIRVConversionOptions &getOptions() const { return options; }
 
   /// Checks if the SPIR-V capability inquired is supported.
-  bool allows(spirv::Capability capability);
+  bool allows(spirv::Capability capability) const;
 
 private:
   spirv::TargetEnv targetEnv;
   SPIRVConversionOptions options;
 
   MLIRContext *getContext() const;
 };
 
Value linearizeIndex(ValueRange indices, ArrayRef<int64_t> strides,
                      OpBuilder &builder);
 
 /// Performs the index computation to get to the element at `indices` of the
 /// memory pointed to by `basePtr`, using the layout map of `baseType`.
 /// Returns null if index computation cannot be performed.
 
 // TODO: This method assumes that the `baseType` is a MemRefType with AffineMap
 // that has static strides. Extend to handle dynamic strides.
-Value getElementPtr(SPIRVTypeConverter &typeConverter, MemRefType baseType,
-                    Value basePtr, ValueRange indices, Location loc,
-                    OpBuilder &builder);
+Value getElementPtr(const SPIRVTypeConverter &typeConverter,
+                    MemRefType baseType, Value basePtr, ValueRange indices,
+                    Location loc, OpBuilder &builder);
 
 // GetElementPtr implementation for Kernel/OpenCL flavored SPIR-V.
-Value getOpenCLElementPtr(SPIRVTypeConverter &typeConverter,
+Value getOpenCLElementPtr(const SPIRVTypeConverter &typeConverter,
                           MemRefType baseType, Value basePtr,
                           ValueRange indices, Location loc, OpBuilder &builder);
 
 // GetElementPtr implementation for Vulkan/Shader flavored SPIR-V.
-Value getVulkanElementPtr(SPIRVTypeConverter &typeConverter,
+Value getVulkanElementPtr(const SPIRVTypeConverter &typeConverter,
                           MemRefType baseType, Value basePtr,
                           ValueRange indices, Location loc, OpBuilder &builder);
 
 } // namespace spirv
 } // namespace mlir
 
 #endif // MLIR_DIALECT_SPIRV_TRANSFORMS_SPIRVCONVERSION_H

mlir/include/mlir/Transforms/DialectConversion.h

public:
   }
 
   /// Attempt to match and rewrite the IR root at the specified operation.
   LogicalResult matchAndRewrite(Operation *op,
                                 PatternRewriter &rewriter) const final;
 
   /// 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:
   /// See `RewritePattern::RewritePattern` for information on the other
   /// available constructors.
   using RewritePattern::RewritePattern;
   /// Construct a conversion pattern with the given converter, and forward the
   /// remaining arguments to RewritePattern.
   template <typename... Args>
-  ConversionPattern(TypeConverter &typeConverter, Args &&...args)
+  ConversionPattern(const TypeConverter &typeConverter, Args &&...args)
       : RewritePattern(std::forward<Args>(args)...),
         typeConverter(&typeConverter) {}
 
 protected:
   /// An optional type converter for use by this pattern.
-  TypeConverter *typeConverter = nullptr;
+  const TypeConverter *typeConverter = nullptr;
 
 private:
   using RewritePattern::rewrite;
 };
 
 /// OpConversionPattern is a wrapper around ConversionPattern that allows for
 /// matching and rewriting against an instance of a derived operation class as
 /// opposed to a raw Operation.
 template <typename SourceOp>
 class OpConversionPattern : public ConversionPattern {
 public:
   using OpAdaptor = typename SourceOp::Adaptor;
 
   OpConversionPattern(MLIRContext *context, PatternBenefit benefit = 1)
       : ConversionPattern(SourceOp::getOperationName(), benefit, context) {}
-  OpConversionPattern(TypeConverter &typeConverter, MLIRContext *context,
+  OpConversionPattern(const TypeConverter &typeConverter, MLIRContext *context,
                       PatternBenefit benefit = 1)
       : ConversionPattern(typeConverter, SourceOp::getOperationName(), benefit,
                           context) {}
 
   /// Wrappers around the ConversionPattern methods that pass the derived op
   /// type.
   LogicalResult match(Operation *op) const final {
     return match(cast<SourceOp>(op));

 /// allows for matching and rewriting against an instance of an OpInterface
 /// class as opposed to a raw Operation.
 template <typename SourceOp>
 class OpInterfaceConversionPattern : public ConversionPattern {
 public:
   OpInterfaceConversionPattern(MLIRContext *context, PatternBenefit benefit = 1)
       : ConversionPattern(Pattern::MatchInterfaceOpTypeTag(),
                           SourceOp::getInterfaceID(), benefit, context) {}
-  OpInterfaceConversionPattern(TypeConverter &typeConverter,
+  OpInterfaceConversionPattern(const TypeConverter &typeConverter,
                                MLIRContext *context, PatternBenefit benefit = 1)
       : ConversionPattern(typeConverter, Pattern::MatchInterfaceOpTypeTag(),
                           SourceOp::getInterfaceID(), benefit, context) {}
 
   /// Wrappers around the ConversionPattern methods that pass the derived op
   /// type.
   void rewrite(Operation *op, ArrayRef<Value> operands,
                ConversionPatternRewriter &rewriter) const final {
private:
   using ConversionPattern::matchAndRewrite;
 };
 
 /// Add a pattern to the given pattern list to convert the signature of a
 /// FunctionOpInterface op with the given type converter. This only supports
 /// ops which use FunctionType to represent their type.
 void populateFunctionOpInterfaceTypeConversionPattern(
     StringRef functionLikeOpName, RewritePatternSet &patterns,
-    TypeConverter &converter);
+    const TypeConverter &converter);
 
 template <typename FuncOpT>
 void populateFunctionOpInterfaceTypeConversionPattern(
-    RewritePatternSet &patterns, TypeConverter &converter) {
+    RewritePatternSet &patterns, const TypeConverter &converter) {
   populateFunctionOpInterfaceTypeConversionPattern(FuncOpT::getOperationName(),
                                                    patterns, converter);
 }
 
 void populateAnyFunctionOpInterfaceTypeConversionPattern(
-    RewritePatternSet &patterns, TypeConverter &converter);
+    RewritePatternSet &patterns, const TypeConverter &converter);
 
 //===----------------------------------------------------------------------===//
 // Conversion PatternRewriter
 //===----------------------------------------------------------------------===//
 
 namespace detail {
 struct ConversionPatternRewriterImpl;
 } // namespace detail
public:
   /// Apply a signature conversion to the entry block of the given region. This
   /// replaces the entry block with a new block containing the updated
   /// signature. The new entry block to the region is returned for convenience.
   ///
   /// If provided, `converter` will be used for any materializations.
   Block *
   applySignatureConversion(Region *region,
                            TypeConverter::SignatureConversion &conversion,
-                           TypeConverter *converter = nullptr);
+                           const TypeConverter *converter = nullptr);
 
   /// Convert the types of block arguments within the given region. This
   /// replaces each block with a new block containing the updated signature. The
   /// entry block may have a special conversion if `entryConversion` is
   /// provided. On success, the new entry block to the region is returned for
   /// convenience. Otherwise, failure is returned.
   FailureOr<Block *> convertRegionTypes(
-      Region *region, TypeConverter &converter,
+      Region *region, const TypeConverter &converter,
       TypeConverter::SignatureConversion *entryConversion = nullptr);
 
   /// Convert the types of block arguments within the given region except for
   /// the entry region. This replaces each non-entry block with a new block
   /// containing the updated signature.
   ///
   /// If special conversion behavior is needed for the non-entry blocks (for
   /// example, we need to convert only a subset of a BB arguments), such
   /// behavior can be specified in blockConversions.
   LogicalResult convertNonEntryRegionTypes(
-      Region *region, TypeConverter &converter,
+      Region *region, const TypeConverter &converter,
       ArrayRef<TypeConverter::SignatureConversion> blockConversions);
 
   /// Replace all the uses of the block argument `from` with value `to`.
   void replaceUsesOfBlockArgument(BlockArgument from, Value to);
 
   /// Return the converted value of 'key' with a type defined by the type
   /// converter of the currently executing pattern. Return nullptr in the case
   /// of failure, the remapped value otherwise.

 // PDL Configuration
 //===----------------------------------------------------------------------===//
 
 /// A PDL configuration that is used to supported dialect conversion
 /// functionality.
 class PDLConversionConfig final
     : public PDLPatternConfigBase<PDLConversionConfig> {
 public:
-  PDLConversionConfig(TypeConverter *converter) : converter(converter) {}
+  PDLConversionConfig(const TypeConverter *converter) : converter(converter) {}
   ~PDLConversionConfig() final = default;
 
   /// Return the type converter used by this configuration, which may be nullptr
   /// if no type conversions are expected.
-  TypeConverter *getTypeConverter() const { return converter; }
+  const TypeConverter *getTypeConverter() const { return converter; }
 
   /// Hooks that are invoked at the beginning and end of a rewrite of a matched
   /// pattern.
   void notifyRewriteBegin(PatternRewriter &rewriter) final;
   void notifyRewriteEnd(PatternRewriter &rewriter) final;
 
 private:
   /// An optional type converter to use for the pattern.
-  TypeConverter *converter;
+  const TypeConverter *converter;
 };
 
 /// Register the dialect conversion PDL functions with the given pattern set.
 void registerConversionPDLFunctions(RewritePatternSet &patterns);
 
 //===----------------------------------------------------------------------===//
 // Op Conversion Entry Points
 //===----------------------------------------------------------------------===//

mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp

static Value createI1Constant(ConversionPatternRewriter &rewriter, Location loc,
   Type llvmI1 = rewriter.getI1Type();
   return rewriter.createOrFold<LLVM::ConstantOp>(loc, llvmI1, value);
 }
 
 namespace {
 /// Define lowering patterns for raw buffer ops
 template <typename GpuOp, typename Intrinsic>
 struct RawBufferOpLowering : public ConvertOpToLLVMPattern<GpuOp> {
-  RawBufferOpLowering(LLVMTypeConverter &converter, Chipset chipset)
+  RawBufferOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
       : ConvertOpToLLVMPattern<GpuOp>(converter), chipset(chipset) {}
 
   Chipset chipset;
   static constexpr uint32_t maxVectorOpWidth = 128;
 
   LogicalResult
   matchAndRewrite(GpuOp gpuOp, typename GpuOp::Adaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
static Value mfmaConcatIfNeeded(ConversionPatternRewriter &rewriter,
   return input;
 }
 
 /// Push an input operand. If it is a float type, nothing to do. If it is
 /// an integer type, then we need to also push its signdness (1 for signed, 0
 /// for unsigned) and we need to pack the input 16xi8 vector into a 4xi32
 /// vector.
 static void wmmaPushInputOperand(ConversionPatternRewriter &rewriter,
-                                 Location loc, TypeConverter *typeConverter,
+                                 Location loc,
+                                 const TypeConverter *typeConverter,
                                  bool isUnsigned, Value llvmInput,
                                  SmallVector<Value, 4> &operands) {
   Type inputType = llvmInput.getType();
   auto vectorType = inputType.dyn_cast<VectorType>();
   Type elemType = vectorType.getElementType();
 
   if (!elemType.isInteger(8)) {
     operands.push_back(llvmInput);

 
 /// Push the output operand. For many cases this is only pushing the output in
 /// the operand list. But when we have f16 -> f16 or bf16 -> bf16 intrinsics,
 /// since the same numbers of VGPRs is used, we need to decide if to store the
 /// result in the upper 16 bits of the VGPRs or in the lower part. To store the
 /// result in the lower 16 bits, set subwordOffset to 1, otherwise result will
 /// be stored it in the upper part
 static void wmmaPushOutputOperand(ConversionPatternRewriter &rewriter,
-                                  Location loc, TypeConverter *typeConverter,
+                                  Location loc,
+                                  const TypeConverter *typeConverter,
                                   Value output, int32_t subwordOffset,
                                   bool clamp, SmallVector<Value, 4> &operands) {
   Type inputType = output.getType();
   auto vectorType = inputType.dyn_cast<VectorType>();
   Type elemType = vectorType.getElementType();
   operands.push_back(output);
   if (elemType.isF16() || elemType.isBF16()) {
     operands.push_back(createI1Constant(rewriter, loc, subwordOffset));
static std::optional<StringRef> wmmaOpToIntrinsic(WMMAOp wmma,
   } else if (elemSourceType.isInteger(8) && elemDestType.isInteger(32)) {
     return ROCDL::wmma_i32_16x16x16_iu8::getOperationName();
   }
   return std::nullopt;
 }
 
 namespace {
 struct MFMAOpLowering : public ConvertOpToLLVMPattern<MFMAOp> {
-  MFMAOpLowering(LLVMTypeConverter &converter, Chipset chipset)
+  MFMAOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
       : ConvertOpToLLVMPattern<MFMAOp>(converter), chipset(chipset) {}
 
   Chipset chipset;
 
   LogicalResult
   matchAndRewrite(MFMAOp op, MFMAOpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     Location loc = op.getLoc();
loweredOp.addOperands(
          createI32Constant(rewriter, loc, getBlgpField)});
     Operation *lowered = rewriter.create(loweredOp);
     rewriter.replaceOp(op, lowered->getResults());
     return success();
   }
 };
 
 struct WMMAOpLowering : public ConvertOpToLLVMPattern<WMMAOp> {
-  WMMAOpLowering(LLVMTypeConverter &converter, Chipset chipset)
+  WMMAOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
       : ConvertOpToLLVMPattern<WMMAOp>(converter), chipset(chipset) {}
 
   Chipset chipset;
 
   LogicalResult
   matchAndRewrite(WMMAOp op, WMMAOpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     Location loc = op.getLoc();

mlir/lib/Conversion/AsyncToLLVM/AsyncToLLVM.cpp

public:
 
   /// Returns whether LLVM opaque pointers should be used instead of typed
   /// pointers.
   bool useOpaquePointers() const { return llvmOpaquePointers; }
 
   /// Creates an LLVM pointer type which may either be a typed pointer or an
   /// opaque pointer, depending on what options the converter was constructed
   /// with.
-  LLVM::LLVMPointerType getPointerType(Type elementType) {
+  LLVM::LLVMPointerType getPointerType(Type elementType) const {
     if (llvmOpaquePointers)
       return LLVM::LLVMPointerType::get(elementType.getContext());
     return LLVM::LLVMPointerType::get(elementType);
   }
 
   static std::optional<Type> convertAsyncTypes(Type type,
                                                bool useOpaquePointers) {
     if (isa<TokenType, GroupType, ValueType>(type))

 /// as type converter. Allows access to it via the 'getTypeConverter'
 /// convenience method.
 template <typename SourceOp>
 class AsyncOpConversionPattern : public OpConversionPattern<SourceOp> {
 
   using Base = OpConversionPattern<SourceOp>;
 
 public:
-  AsyncOpConversionPattern(AsyncRuntimeTypeConverter &typeConverter,
+  AsyncOpConversionPattern(const AsyncRuntimeTypeConverter &typeConverter,
                            MLIRContext *context)
       : Base(typeConverter, context) {}
 
   /// Returns the 'AsyncRuntimeTypeConverter' of the pattern.
-  AsyncRuntimeTypeConverter *getTypeConverter() const {
-    return static_cast<AsyncRuntimeTypeConverter *>(Base::getTypeConverter());
+  const AsyncRuntimeTypeConverter *getTypeConverter() const {
+    return static_cast<const AsyncRuntimeTypeConverter *>(
+        Base::getTypeConverter());
   }
 };
 
 } // namespace
 
 //===----------------------------------------------------------------------===//
 // Convert async.coro.id to @llvm.coro.id intrinsic.
 //===----------------------------------------------------------------------===//

 namespace {
 class RuntimeCreateOpLowering : public ConvertOpToLLVMPattern<RuntimeCreateOp> {
 public:
   using ConvertOpToLLVMPattern::ConvertOpToLLVMPattern;
 
   LogicalResult
   matchAndRewrite(RuntimeCreateOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    TypeConverter *converter = getTypeConverter();
+    const TypeConverter *converter = getTypeConverter();
     Type resultType = op->getResultTypes()[0];
 
     // Tokens creation maps to a simple function call.
     if (isa<TokenType>(resultType)) {
       rewriter.replaceOpWithNewOp<func::CallOp>(
           op, kCreateToken, converter->convertType(resultType));
       return success();
     }

 class RuntimeCreateGroupOpLowering
     : public ConvertOpToLLVMPattern<RuntimeCreateGroupOp> {
 public:
   using ConvertOpToLLVMPattern::ConvertOpToLLVMPattern;
 
   LogicalResult
   matchAndRewrite(RuntimeCreateGroupOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    TypeConverter *converter = getTypeConverter();
+    const TypeConverter *converter = getTypeConverter();
     Type resultType = op.getResult().getType();
 
     rewriter.replaceOpWithNewOp<func::CallOp>(
         op, kCreateGroup, converter->convertType(resultType),
         adaptor.getOperands());
     return success();
   }
 };

 // Async reference counting ops lowering (`async.runtime.add_ref` and
 // `async.runtime.drop_ref` to the corresponding API calls).
 //===----------------------------------------------------------------------===//
 
 namespace {
 template <typename RefCountingOp>
 class RefCountingOpLowering : public OpConversionPattern<RefCountingOp> {
 public:
-  explicit RefCountingOpLowering(TypeConverter &converter, MLIRContext *ctx,
-                                 StringRef apiFunctionName)
+  explicit RefCountingOpLowering(const TypeConverter &converter,
+                                 MLIRContext *ctx, StringRef apiFunctionName)
       : OpConversionPattern<RefCountingOp>(converter, ctx),
         apiFunctionName(apiFunctionName) {}
 
   LogicalResult
   matchAndRewrite(RefCountingOp op, typename RefCountingOp::Adaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     auto count = rewriter.create<arith::ConstantOp>(
         op->getLoc(), rewriter.getI64Type(),
         rewriter.getI64IntegerAttr(op.getCount()));
 
     auto operand = adaptor.getOperand();
     rewriter.replaceOpWithNewOp<func::CallOp>(op, TypeRange(), apiFunctionName,
                                               ValueRange({operand, count}));
 
     return success();
   }
 
 private:
   StringRef apiFunctionName;
 };
 
 class RuntimeAddRefOpLowering : public RefCountingOpLowering<RuntimeAddRefOp> {
 public:
-  explicit RuntimeAddRefOpLowering(TypeConverter &converter, MLIRContext *ctx)
+  explicit RuntimeAddRefOpLowering(const TypeConverter &converter,
+                                   MLIRContext *ctx)
       : RefCountingOpLowering(converter, ctx, kAddRef) {}
 };
 
 class RuntimeDropRefOpLowering
     : public RefCountingOpLowering<RuntimeDropRefOp> {
 public:
-  explicit RuntimeDropRefOpLowering(TypeConverter &converter, MLIRContext *ctx)
+  explicit RuntimeDropRefOpLowering(const TypeConverter &converter,
+                                    MLIRContext *ctx)
       : RefCountingOpLowering(converter, ctx, kDropRef) {}
 };
 } // namespace
 
 //===----------------------------------------------------------------------===//
 // Convert return operations that return async values from async regions.
 //===----------------------------------------------------------------------===//
 

mlir/lib/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.cpp

static std::string generateGlobalMsgSymbolName(ModuleOp moduleOp) {
   int counter = 0;
   while (moduleOp.lookupSymbol(prefix + std::to_string(counter)))
     ++counter;
   return prefix + std::to_string(counter);
 }
 
 /// Generate IR that prints the given string to stderr.
 static void createPrintMsg(OpBuilder &builder, Location loc, ModuleOp moduleOp,
-                           StringRef msg, LLVMTypeConverter &typeConverter) {
+                           StringRef msg,
+                           const LLVMTypeConverter &typeConverter) {
   auto ip = builder.saveInsertionPoint();
   builder.setInsertionPointToStart(moduleOp.getBody());
   MLIRContext *ctx = builder.getContext();
 
   // Create a zero-terminated byte representation and allocate global symbol.
   SmallVector<uint8_t> elementVals;
   elementVals.append(msg.begin(), msg.end());
   elementVals.push_back(0);

mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp

 #define PASS_NAME "convert-func-to-llvm"
 
 static constexpr StringRef varargsAttrName = "func.varargs";
 static constexpr StringRef linkageAttrName = "llvm.linkage";
 static constexpr StringRef barePtrAttrName = "llvm.bareptr";
 
 /// Return `true` if the `op` should use bare pointer calling convention.
 static bool shouldUseBarePtrCallConv(Operation *op,
-                                     LLVMTypeConverter *typeConverter) {
+                                     const LLVMTypeConverter *typeConverter) {
   return (op && op->hasAttr(barePtrAttrName)) ||
          typeConverter->getOptions().useBarePtrCallConv;
 }
 
 /// Only retain those attributes that are not constructed by
 /// `LLVMFuncOp::build`. If `filterArgAttrs` is set, also filter out argument
 /// attributes.
 static void filterFuncAttributes(func::FuncOp func, bool filterArgAndResAttrs,

 /// arguments instead of unpacked arguments. This function can be called from C
 /// by passing a pointer to a C struct corresponding to a memref descriptor.
 /// Similarly, returned memrefs are passed via pointers to a C struct that is
 /// passed as additional argument.
 /// Internally, the auxiliary function unpacks the descriptor into individual
 /// components and forwards them to `newFuncOp` and forwards the results to
 /// the extra arguments.
 static void wrapForExternalCallers(OpBuilder &rewriter, Location loc,
-                                   LLVMTypeConverter &typeConverter,
+                                   const LLVMTypeConverter &typeConverter,
                                    func::FuncOp funcOp,
                                    LLVM::LLVMFuncOp newFuncOp) {
   auto type = funcOp.getFunctionType();
   auto [wrapperFuncType, resultStructType] =
       typeConverter.convertFunctionTypeCWrapper(type);
 
   SmallVector<NamedAttribute, 4> attributes;
   // Only modify the argument and result attributes when the result is now an

 /// `newFuncOp` that allocates a memref descriptor on stack, packs the
 /// individual arguments into this descriptor and passes a pointer to it into
 /// the auxiliary function. If the result of the function cannot be directly
 /// returned, we write it to a special first argument that provides a pointer
 /// to a corresponding struct. This auxiliary external function is now
 /// compatible with functions defined in C using pointers to C structs
 /// corresponding to a memref descriptor.
 static void wrapExternalFunction(OpBuilder &builder, Location loc,
-                                 LLVMTypeConverter &typeConverter,
+                                 const LLVMTypeConverter &typeConverter,
                                  func::FuncOp funcOp,
                                  LLVM::LLVMFuncOp newFuncOp) {
   OpBuilder::InsertionGuard guard(builder);
 
   auto [wrapperType, resultStructType] =
       typeConverter.convertFunctionTypeCWrapper(funcOp.getFunctionType());
   // This conversion can only fail if it could not convert one of the argument
   // types. But since it has been applied to a non-wrapper function before, it
if (resultStructType) {
     builder.create<LLVM::ReturnOp>(loc, call.getResults());
   }
 }
 
 /// Modifies the body of the function to construct the `MemRefDescriptor` from
 /// the bare pointer calling convention lowering of `memref` types.
 static void modifyFuncOpToUseBarePtrCallingConv(
     ConversionPatternRewriter &rewriter, Location loc,
-    LLVMTypeConverter &typeConverter, LLVM::LLVMFuncOp funcOp,
+    const LLVMTypeConverter &typeConverter, LLVM::LLVMFuncOp funcOp,
     TypeRange oldArgTypes) {
   if (funcOp.getBody().empty())
     return;
 
   // Promote bare pointers from memref arguments to memref descriptors at the
   // beginning of the function so that all the memrefs in the function have a
   // uniform representation.
   Block *entryBlock = &funcOp.getBody().front();
convertFuncOpToLLVMFuncOp(func::FuncOp funcOp,
     return newFuncOp;
   }
 };
 
 /// FuncOp legalization pattern that converts MemRef arguments to pointers to
 /// MemRef descriptors (LLVM struct data types) containing all the MemRef type
 /// information.
 struct FuncOpConversion : public FuncOpConversionBase {
-  FuncOpConversion(LLVMTypeConverter &converter)
+  FuncOpConversion(const LLVMTypeConverter &converter)
       : FuncOpConversionBase(converter) {}
 
   LogicalResult
   matchAndRewrite(func::FuncOp funcOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     auto newFuncOp = convertFuncOpToLLVMFuncOp(funcOp, rewriter);
     if (!newFuncOp)
       return failure();

mlir/lib/Conversion/GPUCommon/GPUOpsLowering.h

 
 #include "mlir/Conversion/LLVMCommon/Pattern.h"
 #include "mlir/Dialect/GPU/IR/GPUDialect.h"
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
 
 namespace mlir {
 
 struct GPUFuncOpLowering : ConvertOpToLLVMPattern<gpu::GPUFuncOp> {
-  GPUFuncOpLowering(LLVMTypeConverter &converter, unsigned allocaAddrSpace,
-                    unsigned workgroupAddrSpace, StringAttr kernelAttributeName)
+  GPUFuncOpLowering(const LLVMTypeConverter &converter,
+                    unsigned allocaAddrSpace, unsigned workgroupAddrSpace,
+                    StringAttr kernelAttributeName)
       : ConvertOpToLLVMPattern<gpu::GPUFuncOp>(converter),
         allocaAddrSpace(allocaAddrSpace),
         workgroupAddrSpace(workgroupAddrSpace),
         kernelAttributeName(kernelAttributeName) {}
 
   LogicalResult
   matchAndRewrite(gpu::GPUFuncOp gpuFuncOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override;

 /// The lowering of gpu.printf to a call to an external printf() function
 ///
 /// This pass will add a declaration of printf() to the GPUModule if needed
 /// and seperate out the format strings into global constants. For some
 /// runtimes, such as OpenCL on AMD, this is sufficient setup, as the compiler
 /// will lower printf calls to appropriate device-side code
 struct GPUPrintfOpToLLVMCallLowering
     : public ConvertOpToLLVMPattern<gpu::PrintfOp> {
-  GPUPrintfOpToLLVMCallLowering(LLVMTypeConverter &converter,
+  GPUPrintfOpToLLVMCallLowering(const LLVMTypeConverter &converter,
                                 int addressSpace = 0)
       : ConvertOpToLLVMPattern<gpu::PrintfOp>(converter),
         addressSpace(addressSpace) {}
 
   LogicalResult
   matchAndRewrite(gpu::PrintfOp gpuPrintfOp, gpu::PrintfOpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override;
 
matchAndRewrite(gpu::ReturnOp op, OpAdaptor adaptor,
     return success();
   }
 };
 
 namespace impl {
 /// Unrolls op if it's operating on vectors.
 LogicalResult scalarizeVectorOp(Operation *op, ValueRange operands,
                                 ConversionPatternRewriter &rewriter,
-                                LLVMTypeConverter &converter);
+                                const LLVMTypeConverter &converter);
 } // namespace impl
 
 /// Rewriting that unrolls SourceOp to scalars if it's operating on vectors.
 template <typename SourceOp>
 struct ScalarizeVectorOpLowering : public ConvertOpToLLVMPattern<SourceOp> {
 public:
   using ConvertOpToLLVMPattern<SourceOp>::ConvertOpToLLVMPattern;
 

mlir/lib/Conversion/GPUCommon/GPUOpsLowering.cpp

LogicalResult GPUPrintfOpToVPrintfLowering::matchAndRewrite(
   rewriter.create<LLVM::CallOp>(loc, vprintfDecl, printfArgs);
   rewriter.eraseOp(gpuPrintfOp);
   return success();
 }
 
 /// Unrolls op if it's operating on vectors.
 LogicalResult impl::scalarizeVectorOp(Operation *op, ValueRange operands,
                                       ConversionPatternRewriter &rewriter,
-                                      LLVMTypeConverter &converter) {
+                                      const LLVMTypeConverter &converter) {
   TypeRange operandTypes(operands);
   if (llvm::none_of(operandTypes,
                     [](Type type) { return isa<VectorType>(type); })) {
     return rewriter.notifyMatchFailure(op, "expected vector operand");
   }
   if (op->getNumRegions() != 0 || op->getNumSuccessors() != 0)
     return rewriter.notifyMatchFailure(op, "expected no region/successor");
   if (op->getNumResults() != 1)

mlir/lib/Conversion/GPUCommon/GPUToLLVMConversion.cpp

public:
 
   // Run the dialect converter on the module.
   void runOnOperation() override;
 };
 
 template <typename OpTy>
 class ConvertOpToGpuRuntimeCallPattern : public ConvertOpToLLVMPattern<OpTy> {
 public:
-  explicit ConvertOpToGpuRuntimeCallPattern(LLVMTypeConverter &typeConverter)
+  explicit ConvertOpToGpuRuntimeCallPattern(
+      const LLVMTypeConverter &typeConverter)
       : ConvertOpToLLVMPattern<OpTy>(typeConverter) {}
 
 protected:
   Value getNumElements(ConversionPatternRewriter &rewriter, Location loc,
                        MemRefType type, MemRefDescriptor desc) const {
     Type indexType = ConvertToLLVMPattern::getIndexType();
     return type.hasStaticShape()
                ? ConvertToLLVMPattern::createIndexAttrConstant(
FunctionCallBuilder createSpGEMMGetSizeBuilder = {
        llvmPointerType /*nc*/, llvmPointerType /*void *stream*/}};
 };
 
 /// A rewrite pattern to convert gpu.host_register operations into a GPU runtime
 /// call. Currently it supports CUDA and ROCm (HIP).
 class ConvertHostRegisterOpToGpuRuntimeCallPattern
     : public ConvertOpToGpuRuntimeCallPattern<gpu::HostRegisterOp> {
 public:
-  ConvertHostRegisterOpToGpuRuntimeCallPattern(LLVMTypeConverter &typeConverter)
+  ConvertHostRegisterOpToGpuRuntimeCallPattern(
+      const LLVMTypeConverter &typeConverter)
       : ConvertOpToGpuRuntimeCallPattern<gpu::HostRegisterOp>(typeConverter) {}
 
 private:
   LogicalResult
   matchAndRewrite(gpu::HostRegisterOp hostRegisterOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
 class ConvertHostUnregisterOpToGpuRuntimeCallPattern
     : public ConvertOpToGpuRuntimeCallPattern<gpu::HostUnregisterOp> {
 public:
   ConvertHostUnregisterOpToGpuRuntimeCallPattern(
-      LLVMTypeConverter &typeConverter)
+      const LLVMTypeConverter &typeConverter)
       : ConvertOpToGpuRuntimeCallPattern<gpu::HostUnregisterOp>(typeConverter) {
   }
 
 private:
   LogicalResult
   matchAndRewrite(gpu::HostUnregisterOp hostUnregisterOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
 /// A rewrite pattern to convert gpu.alloc operations into a GPU runtime
 /// call. Currently it supports CUDA and ROCm (HIP).
 class ConvertAllocOpToGpuRuntimeCallPattern
     : public ConvertOpToGpuRuntimeCallPattern<gpu::AllocOp> {
 public:
-  ConvertAllocOpToGpuRuntimeCallPattern(LLVMTypeConverter &typeConverter)
+  ConvertAllocOpToGpuRuntimeCallPattern(const LLVMTypeConverter &typeConverter)
       : ConvertOpToGpuRuntimeCallPattern<gpu::AllocOp>(typeConverter) {}
 
 private:
   LogicalResult
   matchAndRewrite(gpu::AllocOp allocOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
 /// A rewrite pattern to convert gpu.dealloc operations into a GPU runtime
 /// call. Currently it supports CUDA and ROCm (HIP).
 class ConvertDeallocOpToGpuRuntimeCallPattern
     : public ConvertOpToGpuRuntimeCallPattern<gpu::DeallocOp> {
 public:
-  ConvertDeallocOpToGpuRuntimeCallPattern(LLVMTypeConverter &typeConverter)
+  ConvertDeallocOpToGpuRuntimeCallPattern(
+      const LLVMTypeConverter &typeConverter)
       : ConvertOpToGpuRuntimeCallPattern<gpu::DeallocOp>(typeConverter) {}
 
 private:
   LogicalResult
   matchAndRewrite(gpu::DeallocOp deallocOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
 class ConvertAsyncYieldToGpuRuntimeCallPattern
     : public ConvertOpToGpuRuntimeCallPattern<async::YieldOp> {
 public:
-  ConvertAsyncYieldToGpuRuntimeCallPattern(LLVMTypeConverter &typeConverter)
+  ConvertAsyncYieldToGpuRuntimeCallPattern(
+      const LLVMTypeConverter &typeConverter)
       : ConvertOpToGpuRuntimeCallPattern<async::YieldOp>(typeConverter) {}
 
 private:
   LogicalResult
   matchAndRewrite(async::YieldOp yieldOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
 /// A rewrite pattern to convert gpu.wait operations into a GPU runtime
 /// call. Currently it supports CUDA and ROCm (HIP).
 class ConvertWaitOpToGpuRuntimeCallPattern
     : public ConvertOpToGpuRuntimeCallPattern<gpu::WaitOp> {
 public:
-  ConvertWaitOpToGpuRuntimeCallPattern(LLVMTypeConverter &typeConverter)
+  ConvertWaitOpToGpuRuntimeCallPattern(const LLVMTypeConverter &typeConverter)
       : ConvertOpToGpuRuntimeCallPattern<gpu::WaitOp>(typeConverter) {}
 
 private:
   LogicalResult
   matchAndRewrite(gpu::WaitOp waitOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
 /// A rewrite pattern to convert gpu.wait async operations into a GPU runtime
 /// call. Currently it supports CUDA and ROCm (HIP).
 class ConvertWaitAsyncOpToGpuRuntimeCallPattern
     : public ConvertOpToGpuRuntimeCallPattern<gpu::WaitOp> {
 public:
-  ConvertWaitAsyncOpToGpuRuntimeCallPattern(LLVMTypeConverter &typeConverter)
+  ConvertWaitAsyncOpToGpuRuntimeCallPattern(
+      const LLVMTypeConverter &typeConverter)
       : ConvertOpToGpuRuntimeCallPattern<gpu::WaitOp>(typeConverter) {}
 
 private:
   LogicalResult
   matchAndRewrite(gpu::WaitOp waitOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
 /// A rewrite patter to convert gpu.launch_func operations into a sequence of
 /// GPU runtime calls. Currently it supports CUDA and ROCm (HIP).
 ///
 /// In essence, a gpu.launch_func operations gets compiled into the following
 /// sequence of runtime calls:
 ///
 /// * moduleLoad        -- loads the module given the cubin / hsaco data
 /// * moduleGetFunction -- gets a handle to the actual kernel function
 /// * getStreamHelper   -- initializes a new compute stream on GPU
 /// * launchKernel      -- launches the kernel on a stream
 /// * streamSynchronize -- waits for operations on the stream to finish
 ///
 /// Intermediate data structures are allocated on the stack.
 class ConvertLaunchFuncOpToGpuRuntimeCallPattern
     : public ConvertOpToGpuRuntimeCallPattern<gpu::LaunchFuncOp> {
 public:
-  ConvertLaunchFuncOpToGpuRuntimeCallPattern(LLVMTypeConverter &typeConverter,
-                                             StringRef gpuBinaryAnnotation,
-                                             bool kernelBarePtrCallConv)
+  ConvertLaunchFuncOpToGpuRuntimeCallPattern(
+      const LLVMTypeConverter &typeConverter, StringRef gpuBinaryAnnotation,
       : ConvertOpToGpuRuntimeCallPattern<gpu::LaunchFuncOp>(typeConverter),
         gpuBinaryAnnotation(gpuBinaryAnnotation),
         kernelBarePtrCallConv(kernelBarePtrCallConv) {}
 
 private:
   Value generateParamsArray(gpu::LaunchFuncOp launchOp, OpAdaptor adaptor,
                             OpBuilder &builder) const;
   Value generateKernelNameConstant(StringRef moduleName, StringRef name,
class EraseGpuModuleOpPattern : public OpRewritePattern<gpu::GPUModuleOp> {
   }
 };
 
 /// A rewrite pattern to convert gpu.memcpy operations into a GPU runtime
 /// call. Currently it supports CUDA and ROCm (HIP).
 class ConvertMemcpyOpToGpuRuntimeCallPattern
     : public ConvertOpToGpuRuntimeCallPattern<gpu::MemcpyOp> {
 public:
-  ConvertMemcpyOpToGpuRuntimeCallPattern(LLVMTypeConverter &typeConverter)
+  ConvertMemcpyOpToGpuRuntimeCallPattern(const LLVMTypeConverter &typeConverter)
       : ConvertOpToGpuRuntimeCallPattern<gpu::MemcpyOp>(typeConverter) {}
 
 private:
   LogicalResult
   matchAndRewrite(gpu::MemcpyOp memcpyOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
 /// A rewrite pattern to convert gpu.memset operations into a GPU runtime
 /// call. Currently it supports CUDA and ROCm (HIP).
 class ConvertMemsetOpToGpuRuntimeCallPattern
     : public ConvertOpToGpuRuntimeCallPattern<gpu::MemsetOp> {
 public:
-  ConvertMemsetOpToGpuRuntimeCallPattern(LLVMTypeConverter &typeConverter)
+  ConvertMemsetOpToGpuRuntimeCallPattern(const LLVMTypeConverter &typeConverter)
       : ConvertOpToGpuRuntimeCallPattern<gpu::MemsetOp>(typeConverter) {}
 
 private:
   LogicalResult
   matchAndRewrite(gpu::MemsetOp memsetOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
 /// A rewrite pattern to convert gpu.set_default_device to a GPU runtime call.
 /// Currently supports CUDA and ROCm (HIP)
 class ConvertSetDefaultDeviceOpToGpuRuntimeCallPattern
     : public ConvertOpToGpuRuntimeCallPattern<gpu::SetDefaultDeviceOp> {
 public:
   ConvertSetDefaultDeviceOpToGpuRuntimeCallPattern(
-      LLVMTypeConverter &typeConverter)
+      const LLVMTypeConverter &typeConverter)
       : ConvertOpToGpuRuntimeCallPattern<gpu::SetDefaultDeviceOp>(
             typeConverter) {}
 
   LogicalResult
   matchAndRewrite(gpu::SetDefaultDeviceOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
 /// Generic rewriting rule for operation on sparse matrices.
 /// Currently supports CUDA (by means of cuSparse and cuSparseLt).
 #define DECLARE_CONVERT_OP_TO_GPU_RUNTIME_CALL_PATTERN(op_name)                \
   class Convert##op_name##ToGpuRuntimeCallPattern                              \
       : public ConvertOpToGpuRuntimeCallPattern<gpu::op_name> {                \
   public:                                                                      \
     Convert##op_name##ToGpuRuntimeCallPattern(                                 \
-        LLVMTypeConverter &typeConverter)                                      \
+        const LLVMTypeConverter &typeConverter)                                \
         : ConvertOpToGpuRuntimeCallPattern<gpu::op_name>(typeConverter) {}     \
                                                                                \
   private:                                                                     \
     LogicalResult                                                              \
     matchAndRewrite(gpu::op_name op, OpAdaptor adaptor,                        \
                     ConversionPatternRewriter &rewriter) const override;       \
   };
 

 // return %array
 Value ConvertLaunchFuncOpToGpuRuntimeCallPattern::generateParamsArray(
     gpu::LaunchFuncOp launchOp, OpAdaptor adaptor, OpBuilder &builder) const {
   auto loc = launchOp.getLoc();
   auto numKernelOperands = launchOp.getNumKernelOperands();
   SmallVector<Value, 4> arguments;
   if (kernelBarePtrCallConv) {
     // Hack the bare pointer value on just for the argument promotion
-    LLVMTypeConverter *converter = getTypeConverter();
+    const LLVMTypeConverter *converter = getTypeConverter();
     LowerToLLVMOptions options = converter->getOptions();
     LowerToLLVMOptions overrideToMatchKernelOpts = options;
     overrideToMatchKernelOpts.useBarePtrCallConv = true;
-    converter->dangerousSetOptions(overrideToMatchKernelOpts);
-    arguments = converter->promoteOperands(
+    LLVMTypeConverter newConverter = *converter;
+    newConverter.dangerousSetOptions(overrideToMatchKernelOpts);
+    arguments = newConverter.promoteOperands(
         loc, launchOp.getOperands().take_back(numKernelOperands),
         adaptor.getOperands().take_back(numKernelOperands), builder);
-    converter->dangerousSetOptions(options);
   } else {
     arguments = getTypeConverter()->promoteOperands(
         loc, launchOp.getOperands().take_back(numKernelOperands),
         adaptor.getOperands().take_back(numKernelOperands), builder);
   }
 
   auto numArguments = arguments.size();
   SmallVector<Type, 4> argumentTypes;
LogicalResult ConvertLaunchFuncOpToGpuRuntimeCallPattern::matchAndRewrite(
 
   return success();
 }
 
 static Value bitAndAddrspaceCast(Location loc,
                                  ConversionPatternRewriter &rewriter,
                                  LLVM::LLVMPointerType destinationType,
                                  Value sourcePtr,
-                                 LLVMTypeConverter &typeConverter) {
+                                 const LLVMTypeConverter &typeConverter) {
   auto sourceTy = cast<LLVM::LLVMPointerType>(sourcePtr.getType());
   if (destinationType.getAddressSpace() != sourceTy.getAddressSpace())
     sourcePtr = rewriter.create<LLVM::AddrSpaceCastOp>(
         loc,
         typeConverter.getPointerType(sourceTy.getElementType(),
                                      destinationType.getAddressSpace()),
         sourcePtr);
 

mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp

 }
 
 //===----------------------------------------------------------------------===//
 // GPUFuncOp
 //===----------------------------------------------------------------------===//
 
 // Legalizes a GPU function as an entry SPIR-V function.
 static spirv::FuncOp
-lowerAsEntryFunction(gpu::GPUFuncOp funcOp, TypeConverter &typeConverter,
+lowerAsEntryFunction(gpu::GPUFuncOp funcOp, const TypeConverter &typeConverter,
                      ConversionPatternRewriter &rewriter,
                      spirv::EntryPointABIAttr entryPointInfo,
                      ArrayRef<spirv::InterfaceVarABIAttr> argABIInfo) {
   auto fnType = funcOp.getFunctionType();
   if (fnType.getNumResults()) {
     funcOp.emitError("SPIR-V lowering only supports entry functions"
                      "with no return values right now");
     return nullptr;

mlir/lib/Conversion/GPUToSPIRV/WmmaOpsToSPIRV.cpp

matchAndRewrite(gpu::SubgroupMmaLoadMatrixOp subgroupMmaLoadMatrixOp,
                   OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     Location loc = subgroupMmaLoadMatrixOp->getLoc();
     gpu::MMAMatrixType retType =
         cast<gpu::MMAMatrixType>(subgroupMmaLoadMatrixOp.getRes().getType());
     auto memrefType =
         cast<MemRefType>(subgroupMmaLoadMatrixOp.getSrcMemref().getType());
     Value bufferPtr = spirv::getElementPtr(
-        *getTypeConverter<SPIRVTypeConverter>(), memrefType,
+        *getTypeConverter<const SPIRVTypeConverter>(), memrefType,
         adaptor.getSrcMemref(), adaptor.getIndices(), loc, rewriter);
     auto coopType = convertMMAToSPIRVType(retType);
     int64_t stride = subgroupMmaLoadMatrixOp.getLeadDimension().getSExtValue();
     auto i32Type = rewriter.getI32Type();
     auto strideValue = rewriter.create<spirv::ConstantOp>(
         loc, i32Type, IntegerAttr::get(i32Type, stride));
     bool isColMajor = static_cast<bool>(subgroupMmaLoadMatrixOp.getTranspose());
     auto columnMajor = rewriter.create<spirv::ConstantOp>(
struct WmmaStoreOpToSPIRVLowering
   LogicalResult
   matchAndRewrite(gpu::SubgroupMmaStoreMatrixOp subgroupMmaStoreMatrixOp,
                   OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     Location loc = subgroupMmaStoreMatrixOp->getLoc();
     auto memrefType =
         cast<MemRefType>(subgroupMmaStoreMatrixOp.getDstMemref().getType());
     Value bufferPtr = spirv::getElementPtr(
-        *getTypeConverter<SPIRVTypeConverter>(), memrefType,
+        *getTypeConverter<const SPIRVTypeConverter>(), memrefType,
         adaptor.getDstMemref(), adaptor.getIndices(), loc, rewriter);
     int64_t stride = subgroupMmaStoreMatrixOp.getLeadDimension().getSExtValue();
     auto i32Type = rewriter.getI32Type();
     auto strideValue = rewriter.create<spirv::ConstantOp>(
         loc, i32Type, IntegerAttr::get(i32Type, stride));
     bool useColMajor =
         static_cast<bool>(subgroupMmaStoreMatrixOp.getTranspose());
     auto columnMajor = rewriter.create<spirv::ConstantOp>(

mlir/lib/Conversion/LLVMCommon/MemRefBuilder.cpp

MemRefDescriptor MemRefDescriptor::undef(OpBuilder &builder, Location loc,
   return MemRefDescriptor(descriptor);
 }
 
 /// Builds IR creating a MemRef descriptor that represents `type` and
 /// populates it with static shape and stride information extracted from the
 /// type.
 MemRefDescriptor
 MemRefDescriptor::fromStaticShape(OpBuilder &builder, Location loc,
-                                  LLVMTypeConverter &typeConverter,
+                                  const LLVMTypeConverter &typeConverter,
                                   MemRefType type, Value memory) {
   return fromStaticShape(builder, loc, typeConverter, type, memory, memory);
 }
 
 MemRefDescriptor MemRefDescriptor::fromStaticShape(
-    OpBuilder &builder, Location loc, LLVMTypeConverter &typeConverter,
+    OpBuilder &builder, Location loc, const LLVMTypeConverter &typeConverter,
     MemRefType type, Value memory, Value alignedMemory) {
   assert(type.hasStaticShape() && "unexpected dynamic shape");
 
   // Extract all strides and offsets and verify they are static.
   auto [strides, offset] = getStridesAndOffset(type);
   assert(!ShapedType::isDynamic(offset) &&
          "expected static offset");
   assert(!llvm::any_of(strides, ShapedType::isDynamic) &&

 
 LLVM::LLVMPointerType MemRefDescriptor::getElementPtrType() {
   return cast<LLVM::LLVMPointerType>(
       cast<LLVM::LLVMStructType>(value.getType())
           .getBody()[kAlignedPtrPosInMemRefDescriptor]);
 }
 
 Value MemRefDescriptor::bufferPtr(OpBuilder &builder, Location loc,
-                                  LLVMTypeConverter &converter,
+                                  const LLVMTypeConverter &converter,
                                   MemRefType type) {
   // When we convert to LLVM, the input memref must have been normalized
   // beforehand. Hence, this call is guaranteed to work.
   auto [strides, offsetCst] = getStridesAndOffset(type);
 
   Value ptr = alignedPtr(builder, loc);
   // For zero offsets, we already have the base pointer.
   if (offsetCst == 0)

 /// composing that descriptor, in the following order:
 /// - allocated pointer;
 /// - aligned pointer;
 /// - offset;
 /// - <rank> sizes;
 /// - <rank> shapes;
 /// where <rank> is the MemRef rank as provided in `type`.
 Value MemRefDescriptor::pack(OpBuilder &builder, Location loc,
-                             LLVMTypeConverter &converter, MemRefType type,
-                             ValueRange values) {
+                             const LLVMTypeConverter &converter,
+                             MemRefType type, ValueRange values) {
   Type llvmType = converter.convertType(type);
   auto d = MemRefDescriptor::undef(builder, loc, llvmType);
 
   d.setAllocatedPtr(builder, loc, values[kAllocatedPtrPosInMemRefDescriptor]);
   d.setAlignedPtr(builder, loc, values[kAlignedPtrPosInMemRefDescriptor]);
   d.setOffset(builder, loc, values[kOffsetPosInMemRefDescriptor]);
 
   int64_t rank = type.getRank();
void UnrankedMemRefDescriptor::setMemRefDescPtr(OpBuilder &builder,
   setPtr(builder, loc, kPtrInUnrankedMemRefDescriptor, v);
 }
 
 /// Builds IR populating an unranked MemRef descriptor structure from a list
 /// of individual constituent values in the following order:
 /// - rank of the memref;
 /// - pointer to the memref descriptor.
 Value UnrankedMemRefDescriptor::pack(OpBuilder &builder, Location loc,
-                                     LLVMTypeConverter &converter,
+                                     const LLVMTypeConverter &converter,
                                      UnrankedMemRefType type,
                                      ValueRange values) {
   Type llvmType = converter.convertType(type);
   auto d = UnrankedMemRefDescriptor::undef(builder, loc, llvmType);
 
   d.setRank(builder, loc, values[kRankInUnrankedMemRefDescriptor]);
   d.setMemRefDescPtr(builder, loc, values[kPtrInUnrankedMemRefDescriptor]);
   return d;
 }
 
 /// Builds IR extracting individual elements that compose an unranked memref
 /// descriptor and returns them as `results` list.
 void UnrankedMemRefDescriptor::unpack(OpBuilder &builder, Location loc,
                                       Value packed,
                                       SmallVectorImpl<Value> &results) {
   UnrankedMemRefDescriptor d(packed);
   results.reserve(results.size() + 2);
   results.push_back(d.rank(builder, loc));
   results.push_back(d.memRefDescPtr(builder, loc));
 }
 
 void UnrankedMemRefDescriptor::computeSizes(
-    OpBuilder &builder, Location loc, LLVMTypeConverter &typeConverter,
+    OpBuilder &builder, Location loc, const LLVMTypeConverter &typeConverter,
     ArrayRef<UnrankedMemRefDescriptor> values, ArrayRef<unsigned> addressSpaces,
     SmallVectorImpl<Value> &sizes) {
   if (values.empty())
     return;
   assert(values.size() == addressSpaces.size() &&
          "must provide address space for each descriptor");
   // Cache the index type.
   Type indexType = typeConverter.getIndexType();
castToElemPtrPtr(OpBuilder &builder, Location loc, Value memRefDescPtr,
   } else {
     elemPtrPtrType = LLVM::LLVMPointerType::get(elemPtrType);
     elementPtrPtr =
         builder.create<LLVM::BitcastOp>(loc, elemPtrPtrType, memRefDescPtr);
   }
   return {elementPtrPtr, elemPtrPtrType};
 }
 
-Value UnrankedMemRefDescriptor::alignedPtr(OpBuilder &builder, Location loc,
-                                           LLVMTypeConverter &typeConverter,
-                                           Value memRefDescPtr,
-                                           LLVM::LLVMPointerType elemPtrType) {
+Value UnrankedMemRefDescriptor::alignedPtr(
+    OpBuilder &builder, Location loc, const LLVMTypeConverter &typeConverter,
+    Value memRefDescPtr, LLVM::LLVMPointerType elemPtrType) {
   auto [elementPtrPtr, elemPtrPtrType] =
       castToElemPtrPtr(builder, loc, memRefDescPtr, elemPtrType);
 
   Value alignedGep =
       builder.create<LLVM::GEPOp>(loc, elemPtrPtrType, elemPtrType,
                                   elementPtrPtr, ArrayRef<LLVM::GEPArg>{1});
   return builder.create<LLVM::LoadOp>(loc, elemPtrType, alignedGep);
 }
 
-void UnrankedMemRefDescriptor::setAlignedPtr(OpBuilder &builder, Location loc,
-                                             LLVMTypeConverter &typeConverter,
-                                             Value memRefDescPtr,
-                                             LLVM::LLVMPointerType elemPtrType,
-                                             Value alignedPtr) {
+void UnrankedMemRefDescriptor::setAlignedPtr(
+    OpBuilder &builder, Location loc, const LLVMTypeConverter &typeConverter,
+    Value memRefDescPtr, LLVM::LLVMPointerType elemPtrType, Value alignedPtr) {
   auto [elementPtrPtr, elemPtrPtrType] =
       castToElemPtrPtr(builder, loc, memRefDescPtr, elemPtrType);
 
   Value alignedGep =
       builder.create<LLVM::GEPOp>(loc, elemPtrPtrType, elemPtrType,
                                   elementPtrPtr, ArrayRef<LLVM::GEPArg>{1});
   builder.create<LLVM::StoreOp>(loc, alignedPtr, alignedGep);
 }
 
 Value UnrankedMemRefDescriptor::offsetBasePtr(
-    OpBuilder &builder, Location loc, LLVMTypeConverter &typeConverter,
+    OpBuilder &builder, Location loc, const LLVMTypeConverter &typeConverter,
     Value memRefDescPtr, LLVM::LLVMPointerType elemPtrType) {
   auto [elementPtrPtr, elemPtrPtrType] =
       castToElemPtrPtr(builder, loc, memRefDescPtr, elemPtrType);
 
   Value offsetGep =
       builder.create<LLVM::GEPOp>(loc, elemPtrPtrType, elemPtrType,
                                   elementPtrPtr, ArrayRef<LLVM::GEPArg>{2});
 
   if (!elemPtrType.isOpaque()) {
     offsetGep = builder.create<LLVM::BitcastOp>(
         loc, LLVM::LLVMPointerType::get(typeConverter.getIndexType()),
         offsetGep);
   }
   return offsetGep;
 }
 
 Value UnrankedMemRefDescriptor::offset(OpBuilder &builder, Location loc,
-                                       LLVMTypeConverter &typeConverter,
+                                       const LLVMTypeConverter &typeConverter,
                                        Value memRefDescPtr,
                                        LLVM::LLVMPointerType elemPtrType) {
   Value offsetPtr =
       offsetBasePtr(builder, loc, typeConverter, memRefDescPtr, elemPtrType);
   return builder.create<LLVM::LoadOp>(loc, typeConverter.getIndexType(),
                                       offsetPtr);
 }
 
 void UnrankedMemRefDescriptor::setOffset(OpBuilder &builder, Location loc,
-                                         LLVMTypeConverter &typeConverter,
+                                         const LLVMTypeConverter &typeConverter,
                                          Value memRefDescPtr,
                                          LLVM::LLVMPointerType elemPtrType,
                                          Value offset) {
   Value offsetPtr =
       offsetBasePtr(builder, loc, typeConverter, memRefDescPtr, elemPtrType);
   builder.create<LLVM::StoreOp>(loc, offset, offsetPtr);
 }
 
-Value UnrankedMemRefDescriptor::sizeBasePtr(OpBuilder &builder, Location loc,
-                                            LLVMTypeConverter &typeConverter,
-                                            Value memRefDescPtr,
-                                            LLVM::LLVMPointerType elemPtrType) {
+Value UnrankedMemRefDescriptor::sizeBasePtr(
+    OpBuilder &builder, Location loc, const LLVMTypeConverter &typeConverter,
+    Value memRefDescPtr, LLVM::LLVMPointerType elemPtrType) {
   Type indexTy = typeConverter.getIndexType();
   Type structTy = LLVM::LLVMStructType::getLiteral(
       indexTy.getContext(), {elemPtrType, elemPtrType, indexTy, indexTy});
   Value structPtr;
   if (elemPtrType.isOpaque()) {
     structPtr = memRefDescPtr;
   } else {
     Type structPtrTy = LLVM::LLVMPointerType::get(structTy);
     structPtr =
         builder.create<LLVM::BitcastOp>(loc, structPtrTy, memRefDescPtr);
   }
 
   auto resultType = elemPtrType.isOpaque()
                         ? LLVM::LLVMPointerType::get(indexTy.getContext())
                         : LLVM::LLVMPointerType::get(indexTy);
   return builder.create<LLVM::GEPOp>(loc, resultType, structTy, structPtr,
                                      ArrayRef<LLVM::GEPArg>{0, 3});
 }
 
 Value UnrankedMemRefDescriptor::size(OpBuilder &builder, Location loc,
-                                     LLVMTypeConverter &typeConverter,
+                                     const LLVMTypeConverter &typeConverter,
                                      Value sizeBasePtr, Value index) {
 
   Type indexTy = typeConverter.getIndexType();
   Type indexPtrTy = typeConverter.getPointerType(indexTy);
 
   Value sizeStoreGep =
       builder.create<LLVM::GEPOp>(loc, indexPtrTy, indexTy, sizeBasePtr, index);
   return builder.create<LLVM::LoadOp>(loc, indexTy, sizeStoreGep);
 }
 
 void UnrankedMemRefDescriptor::setSize(OpBuilder &builder, Location loc,
-                                       LLVMTypeConverter &typeConverter,
+                                       const LLVMTypeConverter &typeConverter,
                                        Value sizeBasePtr, Value index,
                                        Value size) {
   Type indexTy = typeConverter.getIndexType();
   Type indexPtrTy = typeConverter.getPointerType(indexTy);
 
   Value sizeStoreGep =
       builder.create<LLVM::GEPOp>(loc, indexPtrTy, indexTy, sizeBasePtr, index);
   builder.create<LLVM::StoreOp>(loc, size, sizeStoreGep);
 }
 
-Value UnrankedMemRefDescriptor::strideBasePtr(OpBuilder &builder, Location loc,
-                                              LLVMTypeConverter &typeConverter,
-                                              Value sizeBasePtr, Value rank) {
+Value UnrankedMemRefDescriptor::strideBasePtr(
+    OpBuilder &builder, Location loc, const LLVMTypeConverter &typeConverter,
   Type indexTy = typeConverter.getIndexType();
   Type indexPtrTy = typeConverter.getPointerType(indexTy);
 
   return builder.create<LLVM::GEPOp>(loc, indexPtrTy, indexTy, sizeBasePtr,
                                      rank);
 }
 
 Value UnrankedMemRefDescriptor::stride(OpBuilder &builder, Location loc,
-                                       LLVMTypeConverter &typeConverter,
+                                       const LLVMTypeConverter &typeConverter,
                                        Value strideBasePtr, Value index,
                                        Value stride) {
   Type indexTy = typeConverter.getIndexType();
   Type indexPtrTy = typeConverter.getPointerType(indexTy);
 
   Value strideStoreGep = builder.create<LLVM::GEPOp>(loc, indexPtrTy, indexTy,
                                                      strideBasePtr, index);
   return builder.create<LLVM::LoadOp>(loc, indexTy, strideStoreGep);
 }
 
 void UnrankedMemRefDescriptor::setStride(OpBuilder &builder, Location loc,
-                                         LLVMTypeConverter &typeConverter,
+                                         const LLVMTypeConverter &typeConverter,
                                          Value strideBasePtr, Value index,
                                          Value stride) {
   Type indexTy = typeConverter.getIndexType();
   Type indexPtrTy = typeConverter.getPointerType(indexTy);
 
   Value strideStoreGep = builder.create<LLVM::GEPOp>(loc, indexPtrTy, indexTy,
                                                      strideBasePtr, index);
   builder.create<LLVM::StoreOp>(loc, stride, strideStoreGep);
 }

mlir/lib/Conversion/LLVMCommon/Pattern.cpp

 #include "mlir/IR/BuiltinAttributes.h"
 
 using namespace mlir;
 
 //===----------------------------------------------------------------------===//
 // ConvertToLLVMPattern
 //===----------------------------------------------------------------------===//
 
-ConvertToLLVMPattern::ConvertToLLVMPattern(StringRef rootOpName,
-                                           MLIRContext *context,
-                                           LLVMTypeConverter &typeConverter,
-                                           PatternBenefit benefit)
+ConvertToLLVMPattern::ConvertToLLVMPattern(
+    StringRef rootOpName, MLIRContext *context,
+    const LLVMTypeConverter &typeConverter, PatternBenefit benefit)
     : ConversionPattern(typeConverter, rootOpName, benefit, context) {}
 
-LLVMTypeConverter *ConvertToLLVMPattern::getTypeConverter() const {
-  return static_cast<LLVMTypeConverter *>(
+const LLVMTypeConverter *ConvertToLLVMPattern::getTypeConverter() const {
+  return static_cast<const LLVMTypeConverter *>(
       ConversionPattern::getTypeConverter());
 }
 
 LLVM::LLVMDialect &ConvertToLLVMPattern::getDialect() const {
   return *getTypeConverter()->getDialect();
 }
 
 Type ConvertToLLVMPattern::getIndexType() const {

 }
 
 //===----------------------------------------------------------------------===//
 // Detail methods
 //===----------------------------------------------------------------------===//
 
 /// Replaces the given operation "op" with a new operation of type "targetOp"
 /// and given operands.
-LogicalResult LLVM::detail::oneToOneRewrite(
-    Operation *op, StringRef targetOp, ValueRange operands,
-    ArrayRef<NamedAttribute> targetAttrs, LLVMTypeConverter &typeConverter,
+LogicalResult
+LLVM::detail::oneToOneRewrite(Operation *op, StringRef targetOp,
+                              ValueRange operands,
+                              ArrayRef<NamedAttribute> targetAttrs,
+                              const LLVMTypeConverter &typeConverter,
                               ConversionPatternRewriter &rewriter) {
   unsigned numResults = op->getNumResults();
 
   SmallVector<Type> resultTypes;
   if (numResults != 0) {
     resultTypes.push_back(
         typeConverter.packOperationResults(op->getResultTypes()));
     if (!resultTypes.back())
       return failure();

mlir/lib/Conversion/LLVMCommon/TypeConverter.cpp

LLVMTypeConverter::LLVMTypeConverter(MLIRContext *ctx,
   });
 
   // Integer memory spaces map to themselves.
   addTypeAttributeConversion(
       [](BaseMemRefType memref, IntegerAttr addrspace) { return addrspace; });
 }
 
 /// Returns the MLIR context.
-MLIRContext &LLVMTypeConverter::getContext() {
+MLIRContext &LLVMTypeConverter::getContext() const {
   return *getDialect()->getContext();
 }
 
-Type LLVMTypeConverter::getIndexType() {
+Type LLVMTypeConverter::getIndexType() const {
   return IntegerType::get(&getContext(), getIndexTypeBitwidth());
 }
 
 LLVM::LLVMPointerType
-LLVMTypeConverter::getPointerType(Type elementType, unsigned int addressSpace) {
+LLVMTypeConverter::getPointerType(Type elementType,
+                                  unsigned int addressSpace) const {
   if (useOpaquePointers())
     return LLVM::LLVMPointerType::get(&getContext(), addressSpace);
   return LLVM::LLVMPointerType::get(elementType, addressSpace);
 }
 
-unsigned LLVMTypeConverter::getPointerBitwidth(unsigned addressSpace) {
+unsigned LLVMTypeConverter::getPointerBitwidth(unsigned addressSpace) const {
   return options.dataLayout.getPointerSizeInBits(addressSpace);
 }
 
-Type LLVMTypeConverter::convertIndexType(IndexType type) {
+Type LLVMTypeConverter::convertIndexType(IndexType type) const {
   return getIndexType();
 }
 
-Type LLVMTypeConverter::convertIntegerType(IntegerType type) {
+Type LLVMTypeConverter::convertIntegerType(IntegerType type) const {
   return IntegerType::get(&getContext(), type.getWidth());
 }
 
-Type LLVMTypeConverter::convertFloatType(FloatType type) {
+Type LLVMTypeConverter::convertFloatType(FloatType type) const {
   if (type.isFloat8E5M2() || type.isFloat8E4M3FN() || type.isFloat8E5M2FNUZ() ||
       type.isFloat8E4M3FNUZ() || type.isFloat8E4M3B11FNUZ())
     return IntegerType::get(&getContext(), type.getWidth());
   return type;
 }
 
 // Convert a `ComplexType` to an LLVM type. The result is a complex number
 // struct with entries for the
 //   1. real part and for the
 //   2. imaginary part.
-Type LLVMTypeConverter::convertComplexType(ComplexType type) {
+Type LLVMTypeConverter::convertComplexType(ComplexType type) const {
   auto elementType = convertType(type.getElementType());
   return LLVM::LLVMStructType::getLiteral(&getContext(),
                                           {elementType, elementType});
 }
 
 // Except for signatures, MLIR function types are converted into LLVM
 // pointer-to-function types.
-Type LLVMTypeConverter::convertFunctionType(FunctionType type) {
+Type LLVMTypeConverter::convertFunctionType(FunctionType type) const {
   SignatureConversion conversion(type.getNumInputs());
   Type converted = convertFunctionSignature(
       type, /*isVariadic=*/false, options.useBarePtrCallConv, conversion);
   if (!converted)
     return {};
   return getPointerType(converted);
 }
 
 // Function types are converted to LLVM Function types by recursively converting
 // argument and result types.  If MLIR Function has zero results, the LLVM
 // Function has one VoidType result.  If MLIR Function has more than one result,
 // they are into an LLVM StructType in their order of appearance.
 Type LLVMTypeConverter::convertFunctionSignature(
     FunctionType funcTy, bool isVariadic, bool useBarePtrCallConv,
-    LLVMTypeConverter::SignatureConversion &result) {
+    LLVMTypeConverter::SignatureConversion &result) const {
   // Select the argument converter depending on the calling convention.
   useBarePtrCallConv = useBarePtrCallConv || options.useBarePtrCallConv;
   auto funcArgConverter = useBarePtrCallConv ? barePtrFuncArgTypeConverter
                                              : structFuncArgTypeConverter;
   // Convert argument types one by one and check for errors.
   for (auto [idx, type] : llvm::enumerate(funcTy.getInputs())) {
     SmallVector<Type, 8> converted;
     if (failed(funcArgConverter(*this, type, converted)))
if (!resultType)
     return {};
   return LLVM::LLVMFunctionType::get(resultType, result.getConvertedTypes(),
                                      isVariadic);
 }
 
 /// Converts the function type to a C-compatible format, in particular using
 /// pointers to memref descriptors for arguments.
 std::pair<LLVM::LLVMFunctionType, LLVM::LLVMStructType>
-LLVMTypeConverter::convertFunctionTypeCWrapper(FunctionType type) {
+LLVMTypeConverter::convertFunctionTypeCWrapper(FunctionType type) const {
   SmallVector<Type, 4> inputs;
 
   Type resultType = type.getNumResults() == 0
                         ? LLVM::LLVMVoidType::get(&getContext())
                         : packFunctionResults(type.getResults());
   if (!resultType)
     return {};
 

 ///    Elem *allocatedPtr;
 ///    Elem *alignedPtr;
 ///    Index offset;
 ///    Index sizes[Rank]; // omitted when rank == 0
 ///    Index strides[Rank]; // omitted when rank == 0
 ///  };
 SmallVector<Type, 5>
 LLVMTypeConverter::getMemRefDescriptorFields(MemRefType type,
-                                             bool unpackAggregates) {
+                                             bool unpackAggregates) const {
   if (!isStrided(type)) {
     emitError(
         UnknownLoc::get(type.getContext()),
         "conversion to strided form failed either due to non-strided layout "
         "maps (which should have been normalized away) or other reasons");
     return {};
   }
 
LLVMTypeConverter::getMemRefDescriptorFields(MemRefType type,
 
   if (unpackAggregates)
     results.insert(results.end(), 2 * rank, indexTy);
   else
     results.insert(results.end(), 2, LLVM::LLVMArrayType::get(indexTy, rank));
   return results;
 }
 
-unsigned LLVMTypeConverter::getMemRefDescriptorSize(MemRefType type,
-                                                    const DataLayout &layout) {
+unsigned
+LLVMTypeConverter::getMemRefDescriptorSize(MemRefType type,
+                                           const DataLayout &layout) const {
   // Compute the descriptor size given that of its components indicated above.
   unsigned space = *getMemRefAddressSpace(type);
   return 2 * llvm::divideCeil(getPointerBitwidth(space), 8) +
          (1 + 2 * type.getRank()) * layout.getTypeSize(getIndexType());
 }
 
 /// Converts MemRefType to LLVMType. A MemRefType is converted to a struct that
 /// packs the descriptor fields as defined by `getMemRefDescriptorFields`.
-Type LLVMTypeConverter::convertMemRefType(MemRefType type) {
+Type LLVMTypeConverter::convertMemRefType(MemRefType type) const {
   // When converting a MemRefType to a struct with descriptor fields, do not
   // unpack the `sizes` and `strides` arrays.
   SmallVector<Type, 5> types =
       getMemRefDescriptorFields(type, /*unpackAggregates=*/false);
   if (types.empty())
     return {};
   return LLVM::LLVMStructType::getLiteral(&getContext(), types);
 }
 
 /// Convert an unranked memref type into a list of non-aggregate LLVM IR types
 /// that will form the unranked memref descriptor. In particular, the fields
 /// for an unranked memref descriptor are:
 /// 1. index-typed rank, the dynamic rank of this MemRef
 /// 2. void* ptr, pointer to the static ranked MemRef descriptor. This will be
 ///    stack allocated (alloca) copy of a MemRef descriptor that got casted to
 ///    be unranked.
-SmallVector<Type, 2> LLVMTypeConverter::getUnrankedMemRefDescriptorFields() {
+SmallVector<Type, 2>
+LLVMTypeConverter::getUnrankedMemRefDescriptorFields() const {
   return {getIndexType(), getPointerType(IntegerType::get(&getContext(), 8))};
 }
 
-unsigned
-LLVMTypeConverter::getUnrankedMemRefDescriptorSize(UnrankedMemRefType type,
-                                                   const DataLayout &layout) {
+unsigned LLVMTypeConverter::getUnrankedMemRefDescriptorSize(
+    UnrankedMemRefType type, const DataLayout &layout) const {
   // Compute the descriptor size given that of its components indicated above.
   unsigned space = *getMemRefAddressSpace(type);
   return layout.getTypeSize(getIndexType()) +
          llvm::divideCeil(getPointerBitwidth(space), 8);
 }
 
-Type LLVMTypeConverter::convertUnrankedMemRefType(UnrankedMemRefType type) {
+Type LLVMTypeConverter::convertUnrankedMemRefType(
+    UnrankedMemRefType type) const {
   if (!convertType(type.getElementType()))
     return {};
   return LLVM::LLVMStructType::getLiteral(&getContext(),
                                           getUnrankedMemRefDescriptorFields());
 }
 
 FailureOr<unsigned>
-LLVMTypeConverter::getMemRefAddressSpace(BaseMemRefType type) {
+LLVMTypeConverter::getMemRefAddressSpace(BaseMemRefType type) const {
   if (!type.getMemorySpace()) // Default memory space -> 0.
     return 0;
   std::optional<Attribute> converted =
       convertTypeAttribute(type, type.getMemorySpace());
   if (!converted)
     return failure();
   if (!(*converted)) // Conversion to default is 0.
     return 0;
bool LLVMTypeConverter::canConvertToBarePtr(BaseMemRefType type) {
   for (int64_t stride : strides)
     if (ShapedType::isDynamic(stride))
       return false;
 
   return !ShapedType::isDynamic(offset);
 }
 
 /// Convert a memref type to a bare pointer to the memref element type.
-Type LLVMTypeConverter::convertMemRefToBarePtr(BaseMemRefType type) {
+Type LLVMTypeConverter::convertMemRefToBarePtr(BaseMemRefType type) const {
   if (!canConvertToBarePtr(type))
     return {};
   Type elementType = convertType(type.getElementType());
   if (!elementType)
     return {};
   FailureOr<unsigned> addressSpace = getMemRefAddressSpace(type);
   if (failed(addressSpace))
     return {};
   return getPointerType(elementType, *addressSpace);
 }
 
 /// Convert an n-D vector type to an LLVM vector type:
 ///  * 0-D `vector<T>` are converted to vector<1xT>
 ///  * 1-D `vector<axT>` remains as is while,
 ///  * n>1 `vector<ax...xkxT>` convert via an (n-1)-D array type to
 ///    `!llvm.array<ax...array<jxvector<kxT>>>`.
 /// As LLVM does not support arrays of scalable vectors, it is assumed that
 /// scalable vectors are always 1-D. This condition could be relaxed once the
 /// missing functionality is added in LLVM
-Type LLVMTypeConverter::convertVectorType(VectorType type) {
+Type LLVMTypeConverter::convertVectorType(VectorType type) const {
   auto elementType = convertType(type.getElementType());
   if (!elementType)
     return {};
   if (type.getShape().empty())
     return VectorType::get({1}, elementType);
   Type vectorType = VectorType::get(type.getShape().back(), elementType,
                                     type.getScalableDims().back());
   assert(LLVM::isCompatibleVectorType(vectorType) &&
          "expected vector type compatible with the LLVM dialect");
   assert(
       (!type.isScalable() || (type.getRank() == 1)) &&
       "expected 1-D scalable vector (n-D scalable vectors are not supported)");
   auto shape = type.getShape();
   for (int i = shape.size() - 2; i >= 0; --i)
     vectorType = LLVM::LLVMArrayType::get(vectorType, shape[i]);
   return vectorType;
 }
 
 /// Convert a type in the context of the default or bare pointer calling
 /// convention. Calling convention sensitive types, such as MemRefType and
 /// UnrankedMemRefType, are converted following the specific rules for the
 /// calling convention. Calling convention independent types are converted
 /// following the default LLVM type conversions.
-Type LLVMTypeConverter::convertCallingConventionType(Type type,
-                                                     bool useBarePtrCallConv) {
+Type LLVMTypeConverter::convertCallingConventionType(
+    Type type, bool useBarePtrCallConv) const {
   if (useBarePtrCallConv)
     if (auto memrefTy = dyn_cast<BaseMemRefType>(type))
       return convertMemRefToBarePtr(memrefTy);
 
   return convertType(type);
 }
 
 /// Promote the bare pointers in 'values' that resulted from memrefs to
 /// descriptors. 'stdTypes' holds they types of 'values' before the conversion
 /// to the LLVM-IR dialect (i.e., MemRefType, or any other builtin type).
 void LLVMTypeConverter::promoteBarePtrsToDescriptors(
     ConversionPatternRewriter &rewriter, Location loc, ArrayRef<Type> stdTypes,
-    SmallVectorImpl<Value> &values) {
+    SmallVectorImpl<Value> &values) const {
   assert(stdTypes.size() == values.size() &&
          "The number of types and values doesn't match");
   for (unsigned i = 0, end = values.size(); i < end; ++i)
     if (auto memrefTy = dyn_cast<MemRefType>(stdTypes[i]))
       values[i] = MemRefDescriptor::fromStaticShape(rewriter, loc, *this,
                                                     memrefTy, values[i]);
 }
 
 /// Convert a non-empty list of types of values produced by an operation into an
 /// LLVM-compatible type. In particular, if more than one value is
 /// produced, create a literal structure with elements that correspond to each
 /// of the types converted with `convertType`.
-Type LLVMTypeConverter::packOperationResults(TypeRange types) {
+Type LLVMTypeConverter::packOperationResults(TypeRange types) const {
   assert(!types.empty() && "expected non-empty list of type");
   if (types.size() == 1)
     return convertType(types[0]);
 
   SmallVector<Type> resultTypes;
   resultTypes.reserve(types.size());
   for (Type type : types) {
     Type converted = convertType(type);
     if (!converted || !LLVM::isCompatibleType(converted))
       return {};
     resultTypes.push_back(converted);
   }
 
   return LLVM::LLVMStructType::getLiteral(&getContext(), resultTypes);
 }
 
 /// Convert a non-empty list of types to be returned from a function into an
 /// LLVM-compatible type. In particular, if more than one value is returned,
 /// create an LLVM dialect structure type with elements that correspond to each
 /// of the types converted with `convertCallingConventionType`.
 Type LLVMTypeConverter::packFunctionResults(TypeRange types,
-                                            bool useBarePtrCallConv) {
+                                            bool useBarePtrCallConv) const {
   assert(!types.empty() && "expected non-empty list of type");
 
   useBarePtrCallConv |= options.useBarePtrCallConv;
   if (types.size() == 1)
     return convertCallingConventionType(types.front(), useBarePtrCallConv);
 
   SmallVector<Type> resultTypes;
   resultTypes.reserve(types.size());
   for (auto t : types) {
     auto converted = convertCallingConventionType(t, useBarePtrCallConv);
     if (!converted || !LLVM::isCompatibleType(converted))
       return {};
     resultTypes.push_back(converted);
   }
 
   return LLVM::LLVMStructType::getLiteral(&getContext(), resultTypes);
 }
 
 Value LLVMTypeConverter::promoteOneMemRefDescriptor(Location loc, Value operand,
-                                                    OpBuilder &builder) {
+                                                    OpBuilder &builder) const {
   // Alloca with proper alignment. We do not expect optimizations of this
   // alloca op and so we omit allocating at the entry block.
   auto ptrType = getPointerType(operand.getType());
   Value one = builder.create<LLVM::ConstantOp>(loc, builder.getI64Type(),
                                                builder.getIndexAttr(1));
   Value allocated =
       builder.create<LLVM::AllocaOp>(loc, ptrType, operand.getType(), one);
   // Store into the alloca'ed descriptor.
   builder.create<LLVM::StoreOp>(loc, operand, allocated);
   return allocated;
 }
 
 SmallVector<Value, 4>
 LLVMTypeConverter::promoteOperands(Location loc, ValueRange opOperands,
                                    ValueRange operands, OpBuilder &builder,
-                                   bool useBarePtrCallConv) {
+                                   bool useBarePtrCallConv) const {
   SmallVector<Value, 4> promotedOperands;
   promotedOperands.reserve(operands.size());
   useBarePtrCallConv |= options.useBarePtrCallConv;
   for (auto it : llvm::zip(opOperands, operands)) {
     auto operand = std::get<0>(it);
     auto llvmOperand = std::get<1>(it);
 
     if (useBarePtrCallConv) {
LLVMTypeConverter::promoteOperands(Location loc, ValueRange opOperands,
   }
   return promotedOperands;
 }
 
 /// Callback to convert function argument types. It converts a MemRef function
 /// argument to a list of non-aggregate types containing descriptor
 /// information, and an UnrankedmemRef function argument to a list containing
 /// the rank and a pointer to a descriptor struct.
-LogicalResult mlir::structFuncArgTypeConverter(LLVMTypeConverter &converter,
-                                               Type type,
-                                               SmallVectorImpl<Type> &result) {
+LogicalResult
+mlir::structFuncArgTypeConverter(const LLVMTypeConverter &converter, Type type,
   if (auto memref = dyn_cast<MemRefType>(type)) {
     // In signatures, Memref descriptors are expanded into lists of
     // non-aggregate values.
     auto converted =
         converter.getMemRefDescriptorFields(memref, /*unpackAggregates=*/true);
     if (converted.empty())
       return failure();
     result.append(converted.begin(), converted.end());
mlir::structFuncArgTypeConverter(const LLVMTypeConverter &converter, Type type,
   if (!converted)
     return failure();
   result.push_back(converted);
   return success();
 }
 
 /// Callback to convert function argument types. It converts MemRef function
 /// arguments to bare pointers to the MemRef element type.
-LogicalResult mlir::barePtrFuncArgTypeConverter(LLVMTypeConverter &converter,
-                                                Type type,
-                                                SmallVectorImpl<Type> &result) {
+LogicalResult
+mlir::barePtrFuncArgTypeConverter(const LLVMTypeConverter &converter, Type type,
   auto llvmTy = converter.convertCallingConventionType(
       type, /*useBarePointerCallConv=*/true);
   if (!llvmTy)
     return failure();
 
   result.push_back(llvmTy);
   return success();
 }

mlir/lib/Conversion/LLVMCommon/VectorPattern.cpp

 using namespace mlir;
 
 // For >1-D vector types, extracts the necessary information to iterate over all
 // 1-D subvectors in the underlying llrepresentation of the n-D vector
 // Iterates on the llvm array type until we hit a non-array type (which is
 // asserted to be an llvm vector type).
 LLVM::detail::NDVectorTypeInfo
 LLVM::detail::extractNDVectorTypeInfo(VectorType vectorType,
-                                      LLVMTypeConverter &converter) {
+                                      const LLVMTypeConverter &converter) {
   assert(vectorType.getRank() > 1 && "expected >1D vector type");
   NDVectorTypeInfo info;
   info.llvmNDVectorTy = converter.convertType(vectorType);
   if (!info.llvmNDVectorTy || !LLVM::isCompatibleType(info.llvmNDVectorTy)) {
     info.llvmNDVectorTy = nullptr;
     return info;
   }
   info.arraySizes.reserve(vectorType.getRank() - 1);
for (unsigned linearIndex = 0; linearIndex < ub; ++linearIndex) {
     if (coords.empty())
       break;
     assert(coords.size() == info.arraySizes.size());
     fun(coords);
   }
 }
 
 LogicalResult LLVM::detail::handleMultidimensionalVectors(
-    Operation *op, ValueRange operands, LLVMTypeConverter &typeConverter,
+    Operation *op, ValueRange operands, const LLVMTypeConverter &typeConverter,
     std::function<Value(Type, ValueRange)> createOperand,
     ConversionPatternRewriter &rewriter) {
   auto resultNDVectorType = cast<VectorType>(op->getResult(0).getType());
   auto resultTypeInfo =
       extractNDVectorTypeInfo(resultNDVectorType, typeConverter);
   auto result1DVectorTy = resultTypeInfo.llvm1DVectorTy;
   auto resultNDVectoryTy = resultTypeInfo.llvmNDVectorTy;
   auto loc = op->getLoc();
   Value desc = rewriter.create<LLVM::UndefOp>(loc, resultNDVectoryTy);
   nDVectorIterate(resultTypeInfo, rewriter, [&](ArrayRef<int64_t> position) {
     // For this unrolled `position` corresponding to the `linearIndex`^th
     // element, extract operand vectors
     SmallVector<Value, 4> extractedOperands;
     for (const auto &operand : llvm::enumerate(operands)) {
       extractedOperands.push_back(rewriter.create<LLVM::ExtractValueOp>(
           loc, operand.value(), position));
     }
     Value newVal = createOperand(result1DVectorTy, extractedOperands);
     desc = rewriter.create<LLVM::InsertValueOp>(loc, desc, newVal, position);
   });
   rewriter.replaceOp(op, desc);
   return success();
 }
 
-LogicalResult LLVM::detail::vectorOneToOneRewrite(
-    Operation *op, StringRef targetOp, ValueRange operands,
-    ArrayRef<NamedAttribute> targetAttrs, LLVMTypeConverter &typeConverter,
+LogicalResult
+LLVM::detail::vectorOneToOneRewrite(Operation *op, StringRef targetOp,
+                                    ValueRange operands,
+                                    ArrayRef<NamedAttribute> targetAttrs,
+                                    const LLVMTypeConverter &typeConverter,
                                     ConversionPatternRewriter &rewriter) {
   assert(!operands.empty());
 
   // Cannot convert ops if their operands are not of LLVM type.
   if (!llvm::all_of(operands.getTypes(), isCompatibleType))
     return failure();
 
   auto llvmNDVectorTy = operands[0].getType();
   if (!isa<LLVM::LLVMArrayType>(llvmNDVectorTy))

mlir/lib/Conversion/MemRefToLLVM/AllocLikeConversion.cpp

 
 using namespace mlir;
 
 namespace {
 // TODO: Fix the LLVM utilities for looking up functions to take Operation*
 // with SymbolTable trait instead of ModuleOp and make similar change here. This
 // allows call sites to use getParentWithTrait<OpTrait::SymbolTable> instead
 // of getParentOfType<ModuleOp> to pass down the operation.
-LLVM::LLVMFuncOp getNotalignedAllocFn(LLVMTypeConverter *typeConverter,
+LLVM::LLVMFuncOp getNotalignedAllocFn(const LLVMTypeConverter *typeConverter,
                                       ModuleOp module, Type indexType) {
   bool useGenericFn = typeConverter->getOptions().useGenericFunctions;
 
   if (useGenericFn)
     return LLVM::lookupOrCreateGenericAllocFn(
         module, indexType, typeConverter->useOpaquePointers());
 
   return LLVM::lookupOrCreateMallocFn(module, indexType,
                                       typeConverter->useOpaquePointers());
 }
 
-LLVM::LLVMFuncOp getAlignedAllocFn(LLVMTypeConverter *typeConverter,
+LLVM::LLVMFuncOp getAlignedAllocFn(const LLVMTypeConverter *typeConverter,
                                    ModuleOp module, Type indexType) {
   bool useGenericFn = typeConverter->getOptions().useGenericFunctions;
 
   if (useGenericFn)
     return LLVM::lookupOrCreateGenericAlignedAllocFn(
         module, indexType, typeConverter->useOpaquePointers());
 
   return LLVM::lookupOrCreateAlignedAllocFn(module, indexType,
Value AllocationOpLLVMLowering::createAligned(
   Value bumped = rewriter.create<LLVM::AddOp>(loc, input, bump);
   Value mod = rewriter.create<LLVM::URemOp>(loc, bumped, alignment);
   return rewriter.create<LLVM::SubOp>(loc, bumped, mod);
 }
 
 static Value castAllocFuncResult(ConversionPatternRewriter &rewriter,
                                  Location loc, Value allocatedPtr,
                                  MemRefType memRefType, Type elementPtrType,
-                                 LLVMTypeConverter &typeConverter) {
+                                 const LLVMTypeConverter &typeConverter) {
   auto allocatedPtrTy = cast<LLVM::LLVMPointerType>(allocatedPtr.getType());
   unsigned memrefAddrSpace = *typeConverter.getMemRefAddressSpace(memRefType);
   if (allocatedPtrTy.getAddressSpace() != memrefAddrSpace)
     allocatedPtr = rewriter.create<LLVM::AddrSpaceCastOp>(
         loc,
         typeConverter.getPointerType(allocatedPtrTy.getElementType(),
                                      memrefAddrSpace),
         allocatedPtr);

mlir/lib/Conversion/MemRefToLLVM/MemRefToLLVM.cpp

 using namespace mlir;
 
 namespace {
 
 bool isStaticStrideOrOffset(int64_t strideOrOffset) {
   return !ShapedType::isDynamic(strideOrOffset);
 }
 
-LLVM::LLVMFuncOp getFreeFn(LLVMTypeConverter *typeConverter, ModuleOp module) {
+LLVM::LLVMFuncOp getFreeFn(const LLVMTypeConverter *typeConverter,
+                           ModuleOp module) {
   bool useGenericFn = typeConverter->getOptions().useGenericFunctions;
 
   if (useGenericFn)
     return LLVM::lookupOrCreateGenericFreeFn(
         module, typeConverter->useOpaquePointers());
 
   return LLVM::lookupOrCreateFreeFn(module, typeConverter->useOpaquePointers());
 }
 
 struct AllocOpLowering : public AllocLikeOpLLVMLowering {
-  AllocOpLowering(LLVMTypeConverter &converter)
+  AllocOpLowering(const LLVMTypeConverter &converter)
       : AllocLikeOpLLVMLowering(memref::AllocOp::getOperationName(),
                                 converter) {}
   std::tuple<Value, Value> allocateBuffer(ConversionPatternRewriter &rewriter,
                                           Location loc, Value sizeBytes,
                                           Operation *op) const override {
     return allocateBufferManuallyAlign(
         rewriter, loc, sizeBytes, op,
         getAlignment(rewriter, loc, cast<memref::AllocOp>(op)));
   }
 };
 
 struct AlignedAllocOpLowering : public AllocLikeOpLLVMLowering {
-  AlignedAllocOpLowering(LLVMTypeConverter &converter)
+  AlignedAllocOpLowering(const LLVMTypeConverter &converter)
       : AllocLikeOpLLVMLowering(memref::AllocOp::getOperationName(),
                                 converter) {}
   std::tuple<Value, Value> allocateBuffer(ConversionPatternRewriter &rewriter,
                                           Location loc, Value sizeBytes,
                                           Operation *op) const override {
     Value ptr = allocateBufferAutoAlign(
         rewriter, loc, sizeBytes, op, &defaultLayout,
         alignedAllocationGetAlignment(rewriter, loc, cast<memref::AllocOp>(op),
                                       &defaultLayout));
     return std::make_tuple(ptr, ptr);
   }
 
 private:
   /// Default layout to use in absence of the corresponding analysis.
   DataLayout defaultLayout;
 };
 
 struct AllocaOpLowering : public AllocLikeOpLLVMLowering {
-  AllocaOpLowering(LLVMTypeConverter &converter)
+  AllocaOpLowering(const LLVMTypeConverter &converter)
       : AllocLikeOpLLVMLowering(memref::AllocaOp::getOperationName(),
                                 converter) {
     setRequiresNumElements();
   }
 
   /// Allocates the underlying buffer using the right call. `allocatedBytePtr`
   /// is set to null for stack allocations. `accessAlignment` is set if
   /// alignment is needed post allocation (for eg. in conjunction with malloc).

 
 /// The base class for lowering realloc op, to support the implementation of
 /// realloc via allocation methods that may or may not support alignment.
 /// A derived class should provide an implementation of allocateBuffer using
 /// the underline allocation methods.
 struct ReallocOpLoweringBase : public AllocationOpLLVMLowering {
   using OpAdaptor = typename memref::ReallocOp::Adaptor;
 
-  ReallocOpLoweringBase(LLVMTypeConverter &converter)
+  ReallocOpLoweringBase(const LLVMTypeConverter &converter)
       : AllocationOpLLVMLowering(memref::ReallocOp::getOperationName(),
                                  converter) {}
 
   /// Allocates the new buffer. Returns the allocated pointer and the
   /// aligned pointer.
   virtual std::tuple<Value, Value>
   allocateBuffer(ConversionPatternRewriter &rewriter, Location loc,
                  Value sizeBytes, memref::ReallocOp op) const = 0;
LogicalResult matchAndRewrite(memref::ReallocOp op, OpAdaptor adaptor,
     return success();
   }
 
 private:
   using ConvertToLLVMPattern::matchAndRewrite;
 };
 
 struct ReallocOpLowering : public ReallocOpLoweringBase {
-  ReallocOpLowering(LLVMTypeConverter &converter)
+  ReallocOpLowering(const LLVMTypeConverter &converter)
       : ReallocOpLoweringBase(converter) {}
   std::tuple<Value, Value> allocateBuffer(ConversionPatternRewriter &rewriter,
                                           Location loc, Value sizeBytes,
                                           memref::ReallocOp op) const override {
     return allocateBufferManuallyAlign(rewriter, loc, sizeBytes, op,
                                        getAlignment(rewriter, loc, op));
   }
 };
 
 struct AlignedReallocOpLowering : public ReallocOpLoweringBase {
-  AlignedReallocOpLowering(LLVMTypeConverter &converter)
+  AlignedReallocOpLowering(const LLVMTypeConverter &converter)
       : ReallocOpLoweringBase(converter) {}
   std::tuple<Value, Value> allocateBuffer(ConversionPatternRewriter &rewriter,
                                           Location loc, Value sizeBytes,
                                           memref::ReallocOp op) const override {
     Value ptr = allocateBufferAutoAlign(
         rewriter, loc, sizeBytes, op, &defaultLayout,
         alignedAllocationGetAlignment(rewriter, loc, op, &defaultLayout));
     return std::make_tuple(ptr, ptr);
matchAndRewrite(memref::AllocaScopeOp allocaScopeOp, OpAdaptor adaptor,
     return success();
   }
 };
 
 struct AssumeAlignmentOpLowering
     : public ConvertOpToLLVMPattern<memref::AssumeAlignmentOp> {
   using ConvertOpToLLVMPattern<
       memref::AssumeAlignmentOp>::ConvertOpToLLVMPattern;
-  explicit AssumeAlignmentOpLowering(LLVMTypeConverter &converter)
+  explicit AssumeAlignmentOpLowering(const LLVMTypeConverter &converter)
       : ConvertOpToLLVMPattern<memref::AssumeAlignmentOp>(converter) {}
 
   LogicalResult
   matchAndRewrite(memref::AssumeAlignmentOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     Value memref = adaptor.getMemref();
     unsigned alignment = op.getAlignment();
     auto loc = op.getLoc();

 };
 
 // A `dealloc` is converted into a call to `free` on the underlying data buffer.
 // The memref descriptor being an SSA value, there is no need to clean it up
 // in any way.
 struct DeallocOpLowering : public ConvertOpToLLVMPattern<memref::DeallocOp> {
   using ConvertOpToLLVMPattern<memref::DeallocOp>::ConvertOpToLLVMPattern;
 
-  explicit DeallocOpLowering(LLVMTypeConverter &converter)
+  explicit DeallocOpLowering(const LLVMTypeConverter &converter)
       : ConvertOpToLLVMPattern<memref::DeallocOp>(converter) {}
 
   LogicalResult
   matchAndRewrite(memref::DeallocOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     // Insert the `free` declaration if it is not already present.
     LLVM::LLVMFuncOp freeFunc =
         getFreeFn(getTypeConverter(), op->getParentOfType<ModuleOp>());
matchAndRewrite(memref::GenericAtomicRMWOp atomicOp, OpAdaptor adaptor,
     // The 'result' of the atomic_rmw op is the newly loaded value.
     rewriter.replaceOp(atomicOp, {newLoaded});
 
     return success();
   }
 };
 
 /// Returns the LLVM type of the global variable given the memref type `type`.
-static Type convertGlobalMemrefTypeToLLVM(MemRefType type,
-                                          LLVMTypeConverter &typeConverter) {
+static Type
+convertGlobalMemrefTypeToLLVM(MemRefType type,
+                              const LLVMTypeConverter &typeConverter) {
   // LLVM type for a global memref will be a multi-dimension array. For
   // declarations or uninitialized global memrefs, we can potentially flatten
   // this to a 1D array. However, for memref.global's with an initial value,
   // we do not intend to flatten the ElementsAttribute when going from std ->
   // LLVM dialect, so the LLVM type needs to me a multi-dimension array.
   Type elementType = typeConverter.convertType(type.getElementType());
   Type arrayTy = elementType;
   // Shape has the outermost dim at index 0, so need to walk it backwards
matchAndRewrite(memref::GlobalOp global, OpAdaptor adaptor,
     return success();
   }
 };
 
 /// GetGlobalMemrefOp is lowered into a Memref descriptor with the pointer to
 /// the first element stashed into the descriptor. This reuses
 /// `AllocLikeOpLowering` to reuse the Memref descriptor construction.
 struct GetGlobalMemrefOpLowering : public AllocLikeOpLLVMLowering {
-  GetGlobalMemrefOpLowering(LLVMTypeConverter &converter)
+  GetGlobalMemrefOpLowering(const LLVMTypeConverter &converter)
       : AllocLikeOpLLVMLowering(memref::GetGlobalOp::getOperationName(),
                                 converter) {}
 
   /// Buffer "allocation" for memref.get_global op is getting the address of
   /// the global variable referenced.
   std::tuple<Value, Value> allocateBuffer(ConversionPatternRewriter &rewriter,
                                           Location loc, Value sizeBytes,
                                           Operation *op) const override {
struct MemorySpaceCastOpLowering
   }
 };
 
 /// Extracts allocated, aligned pointers and offset from a ranked or unranked
 /// memref type. In unranked case, the fields are extracted from the underlying
 /// ranked descriptor.
 static void extractPointersAndOffset(Location loc,
                                      ConversionPatternRewriter &rewriter,
-                                     LLVMTypeConverter &typeConverter,
+                                     const LLVMTypeConverter &typeConverter,
                                      Value originalOperand,
                                      Value convertedOperand,
                                      Value *allocatedPtr, Value *alignedPtr,
                                      Value *offset = nullptr) {
   Type operandType = originalOperand.getType();
   if (isa<MemRefType>(operandType)) {
     MemRefDescriptor desc(convertedOperand);
     *allocatedPtr = desc.allocatedPtr(rewriter, loc);

mlir/lib/Conversion/MemRefToSPIRV/MemRefToSPIRV.cpp

 /// Returns an adjusted spirv::AccessChainOp. Based on the
 /// extension/capabilities, certain integer bitwidths `sourceBits` might not be
 /// supported. During conversion if a memref of an unsupported type is used,
 /// load/stores to this memref need to be modified to use a supported higher
 /// bitwidth `targetBits` and extracting the required bits. For an accessing a
 /// 1D array (spirv.array or spirv.rt_array), the last index is modified to load
 /// the bits needed. The extraction of the actual bits needed are handled
 /// separately. Note that this only works for a 1-D tensor.
-static Value adjustAccessChainForBitwidth(SPIRVTypeConverter &typeConverter,
-                                          spirv::AccessChainOp op,
-                                          int sourceBits, int targetBits,
-                                          OpBuilder &builder) {
+static Value
+adjustAccessChainForBitwidth(const SPIRVTypeConverter &typeConverter,
+                             spirv::AccessChainOp op, int sourceBits,
+                             int targetBits, OpBuilder &builder) {
   assert(targetBits % sourceBits == 0);
   const auto loc = op.getLoc();
   IntegerType targetType = builder.getIntegerType(targetBits);
   IntegerAttr attr =
       builder.getIntegerAttr(targetType, targetBits / sourceBits);
   auto idx = builder.create<spirv::ConstantOp>(loc, targetType, attr);
   auto lastDim = op->getOperand(op.getNumOperands() - 1);
   auto indices = llvm::to_vector<4>(op.getIndices());
public:
   using OpConversionPattern::OpConversionPattern;
 
   LogicalResult
   matchAndRewrite(memref::CastOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     Value src = adaptor.getSource();
     Type srcType = src.getType();
 
-    TypeConverter *converter = getTypeConverter();
+    const TypeConverter *converter = getTypeConverter();
     Type dstType = converter->convertType(op.getType());
     if (srcType != dstType)
       return rewriter.notifyMatchFailure(op, [&](Diagnostic &diag) {
         diag << "types doesn't match: " << srcType << " and " << dstType;
       });
 
     rewriter.replaceOp(op, src);
     return success();

 LogicalResult
 IntLoadOpPattern::matchAndRewrite(memref::LoadOp loadOp, OpAdaptor adaptor,
                                   ConversionPatternRewriter &rewriter) const {
   auto loc = loadOp.getLoc();
   auto memrefType = cast<MemRefType>(loadOp.getMemref().getType());
   if (!memrefType.getElementType().isSignlessInteger())
     return failure();
 
-  auto &typeConverter = *getTypeConverter<SPIRVTypeConverter>();
+  const auto &typeConverter = *getTypeConverter<SPIRVTypeConverter>();
   Value accessChain =
       spirv::getElementPtr(typeConverter, memrefType, adaptor.getMemref(),
                            adaptor.getIndices(), loc, rewriter);
 
   if (!accessChain)
     return failure();
 
   int srcBits = memrefType.getElementType().getIntOrFloatBitWidth();
LogicalResult ReinterpretCastPattern::matchAndRewrite(
   Value src = adaptor.getSource();
   auto srcType = dyn_cast<spirv::PointerType>(src.getType());
 
   if (!srcType)
     return rewriter.notifyMatchFailure(op, [&](Diagnostic &diag) {
       diag << "invalid src type " << src.getType();
     });
 
-  TypeConverter *converter = getTypeConverter();
+  const TypeConverter *converter = getTypeConverter();
 
   auto dstType = converter->convertType<spirv::PointerType>(op.getType());
   if (dstType != srcType)
     return rewriter.notifyMatchFailure(op, [&](Diagnostic &diag) {
       diag << "invalid dst type " << op.getType();
     });
 
   OpFoldResult offset =

mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp

MemRefType nvgpu::getMBarrierMemrefType(MLIRContext *context,
   Attribute memorySpace = nvgpu::getMbarrierMemorySpace(context, barrierType);
   MemRefLayoutAttrInterface layout;
   return MemRefType::get({1}, IntegerType::get(context, 64), layout,
                          memorySpace);
 }
 
 /// Returns the base pointer of the mbarrier object.
 static Value getMbarrierPtr(ConversionPatternRewriter &rewriter,
-                            LLVMTypeConverter &typeConverter,
+                            const LLVMTypeConverter &typeConverter,
                             TypedValue<nvgpu::MBarrierType> barrier,
                             Value barrierMemref) {
   MemRefType memrefType =
       nvgpu::getMBarrierMemrefType(rewriter.getContext(), barrier.getType());
   MemRefDescriptor memRefDescriptor(barrierMemref);
   return memRefDescriptor.bufferPtr(rewriter, barrier.getLoc(), typeConverter,
                                     memrefType);
 }

mlir/lib/Conversion/OpenMPToLLVM/OpenMPToLLVM.cpp

 
 template <typename T>
 struct RegionLessOpWithVarOperandsConversion
     : public ConvertOpToLLVMPattern<T> {
   using ConvertOpToLLVMPattern<T>::ConvertOpToLLVMPattern;
   LogicalResult
   matchAndRewrite(T curOp, typename T::Adaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    TypeConverter *converter = ConvertToLLVMPattern::getTypeConverter();
+    const TypeConverter *converter = ConvertToLLVMPattern::getTypeConverter();
     SmallVector<Type> resTypes;
     if (failed(converter->convertTypes(curOp->getResultTypes(), resTypes)))
       return failure();
     SmallVector<Value> convertedOperands;
     assert(curOp.getNumVariableOperands() ==
                curOp.getOperation()->getNumOperands() &&
            "unexpected non-variable operands");
     for (unsigned idx = 0; idx < curOp.getNumVariableOperands(); ++idx) {

 };
 
 template <typename T>
 struct RegionOpWithVarOperandsConversion : public ConvertOpToLLVMPattern<T> {
   using ConvertOpToLLVMPattern<T>::ConvertOpToLLVMPattern;
   LogicalResult
   matchAndRewrite(T curOp, typename T::Adaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    TypeConverter *converter = ConvertToLLVMPattern::getTypeConverter();
+    const TypeConverter *converter = ConvertToLLVMPattern::getTypeConverter();
     SmallVector<Type> resTypes;
     if (failed(converter->convertTypes(curOp->getResultTypes(), resTypes)))
       return failure();
     SmallVector<Value> convertedOperands;
     assert(curOp.getNumVariableOperands() ==
                curOp.getOperation()->getNumOperands() &&
            "unexpected non-variable operands");
     for (unsigned idx = 0; idx < curOp.getNumVariableOperands(); ++idx) {

 };
 
 template <typename T>
 struct RegionLessOpConversion : public ConvertOpToLLVMPattern<T> {
   using ConvertOpToLLVMPattern<T>::ConvertOpToLLVMPattern;
   LogicalResult
   matchAndRewrite(T curOp, typename T::Adaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    TypeConverter *converter = ConvertToLLVMPattern::getTypeConverter();
+    const TypeConverter *converter = ConvertToLLVMPattern::getTypeConverter();
     SmallVector<Type> resTypes;
     if (failed(converter->convertTypes(curOp->getResultTypes(), resTypes)))
       return failure();
 
     rewriter.replaceOpWithNewOp<T>(curOp, resTypes, adaptor.getOperands(),
                                    curOp->getAttrs());
     return success();
   }
 };
 
 struct AtomicReadOpConversion
     : public ConvertOpToLLVMPattern<omp::AtomicReadOp> {
   using ConvertOpToLLVMPattern<omp::AtomicReadOp>::ConvertOpToLLVMPattern;
   LogicalResult
   matchAndRewrite(omp::AtomicReadOp curOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    TypeConverter *converter = ConvertToLLVMPattern::getTypeConverter();
+    const TypeConverter *converter = ConvertToLLVMPattern::getTypeConverter();
     Type curElementType = curOp.getElementType();
     auto newOp = rewriter.create<omp::AtomicReadOp>(
         curOp.getLoc(), TypeRange(), adaptor.getOperands(), curOp->getAttrs());
     TypeAttr typeAttr = TypeAttr::get(converter->convertType(curElementType));
     newOp.setElementTypeAttr(typeAttr);
     rewriter.eraseOp(curOp);
     return success();
   }

mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVM.cpp

 static VectorType reducedVectorTypeBack(VectorType tp) {
   assert((tp.getRank() > 1) && "unlowerable vector type");
   return VectorType::get(tp.getShape().take_back(), tp.getElementType(),
                          tp.getScalableDims().take_back());
 }
 
 // Helper that picks the proper sequence for inserting.
 static Value insertOne(ConversionPatternRewriter &rewriter,
-                       LLVMTypeConverter &typeConverter, Location loc,
+                       const LLVMTypeConverter &typeConverter, Location loc,
                        Value val1, Value val2, Type llvmType, int64_t rank,
                        int64_t pos) {
   assert(rank > 0 && "0-D vector corner case should have been handled already");
   if (rank == 1) {
     auto idxType = rewriter.getIndexType();
     auto constant = rewriter.create<LLVM::ConstantOp>(
         loc, typeConverter.convertType(idxType),
         rewriter.getIntegerAttr(idxType, pos));
     return rewriter.create<LLVM::InsertElementOp>(loc, llvmType, val1, val2,
                                                   constant);
   }
   return rewriter.create<LLVM::InsertValueOp>(loc, val1, val2, pos);
 }
 
 // Helper that picks the proper sequence for extracting.
 static Value extractOne(ConversionPatternRewriter &rewriter,
-                        LLVMTypeConverter &typeConverter, Location loc,
+                        const LLVMTypeConverter &typeConverter, Location loc,
                         Value val, Type llvmType, int64_t rank, int64_t pos) {
   if (rank <= 1) {
     auto idxType = rewriter.getIndexType();
     auto constant = rewriter.create<LLVM::ConstantOp>(
         loc, typeConverter.convertType(idxType),
         rewriter.getIntegerAttr(idxType, pos));
     return rewriter.create<LLVM::ExtractElementOp>(loc, llvmType, val,
                                                    constant);
   }
   return rewriter.create<LLVM::ExtractValueOp>(loc, val, pos);
 }
 
 // Helper that returns data layout alignment of a memref.
-LogicalResult getMemRefAlignment(LLVMTypeConverter &typeConverter,
+LogicalResult getMemRefAlignment(const LLVMTypeConverter &typeConverter,
                                  MemRefType memrefType, unsigned &align) {
   Type elementTy = typeConverter.convertType(memrefType.getElementType());
   if (!elementTy)
     return failure();
 
   // TODO: this should use the MLIR data layout when it becomes available and
   // stop depending on translation.
   llvm::LLVMContext llvmContext;
   align = LLVM::TypeToLLVMIRTranslator(llvmContext)
               .getPreferredAlignment(elementTy, typeConverter.getDataLayout());
   return success();
 }
 
 // Check if the last stride is non-unit or the memory space is not zero.
 static LogicalResult isMemRefTypeSupported(MemRefType memRefType,
-                                           LLVMTypeConverter &converter) {
+                                           const LLVMTypeConverter &converter) {
   if (!isLastMemrefDimUnitStride(memRefType))
     return failure();
   FailureOr<unsigned> addressSpace =
       converter.getMemRefAddressSpace(memRefType);
   if (failed(addressSpace) || *addressSpace != 0)
     return failure();
   return success();
 }
 
 // Add an index vector component to a base pointer.
 static Value getIndexedPtrs(ConversionPatternRewriter &rewriter, Location loc,
-                            LLVMTypeConverter &typeConverter,
+                            const LLVMTypeConverter &typeConverter,
                             MemRefType memRefType, Value llvmMemref, Value base,
                             Value index, uint64_t vLen) {
   assert(succeeded(isMemRefTypeSupported(memRefType, typeConverter)) &&
          "unsupported memref type");
   auto pType = MemRefDescriptor(llvmMemref).getElementPtrType();
   auto ptrsType = LLVM::getFixedVectorType(pType, vLen);
   return rewriter.create<LLVM::GEPOp>(
       loc, ptrsType, typeConverter.convertType(memRefType.getElementType()),
       base, index);
 }
 
 // Casts a strided element pointer to a vector pointer.  The vector pointer
 // will be in the same address space as the incoming memref type.
 static Value castDataPtr(ConversionPatternRewriter &rewriter, Location loc,
                          Value ptr, MemRefType memRefType, Type vt,
-                         LLVMTypeConverter &converter) {
+                         const LLVMTypeConverter &converter) {
   if (converter.useOpaquePointers())
     return ptr;
 
   unsigned addressSpace = *converter.getMemRefAddressSpace(memRefType);
   auto pType = LLVM::LLVMPointerType::get(vt, addressSpace);
   return rewriter.create<LLVM::BitcastOp>(loc, pType, ptr);
 }
 
if (!isa<LLVM::LLVMArrayType>(llvmNDVectorTy)) {
                                   /*vLen=*/vType.getDimSize(0));
       // Replace with the gather intrinsic.
       rewriter.replaceOpWithNewOp<LLVM::masked_gather>(
           gather, typeConverter->convertType(vType), ptrs, adaptor.getMask(),
           adaptor.getPassThru(), rewriter.getI32IntegerAttr(align));
       return success();
     }
 
-    LLVMTypeConverter &typeConverter = *this->getTypeConverter();
+    const LLVMTypeConverter &typeConverter = *this->getTypeConverter();
     auto callback = [align, memRefType, base, ptr, loc, &rewriter,
                      &typeConverter](Type llvm1DVectorTy,
                                      ValueRange vectorOperands) {
       // Resolve address.
       Value ptrs = getIndexedPtrs(
           rewriter, loc, typeConverter, memRefType, base, ptr,
           /*index=*/vectorOperands[0],
           LLVM::getVectorNumElements(llvm1DVectorTy).getFixedValue());
static Value lowerReductionWithStartValue(ConversionPatternRewriter &rewriter,
   return lowerReductionWithStartValue<LLVMFPVPRedIntrinOp, FPReductionNeutral>(
       rewriter, loc, llvmType, vectorOperand, accumulator, mask);
 }
 
 /// Conversion pattern for all vector reductions.
 class VectorReductionOpConversion
     : public ConvertOpToLLVMPattern<vector::ReductionOp> {
 public:
-  explicit VectorReductionOpConversion(LLVMTypeConverter &typeConv,
+  explicit VectorReductionOpConversion(const LLVMTypeConverter &typeConv,
                                        bool reassociateFPRed)
       : ConvertOpToLLVMPattern<vector::ReductionOp>(typeConv),
         reassociateFPReductions(reassociateFPRed) {}
 
   LogicalResult
   matchAndRewrite(vector::ReductionOp reductionOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     auto kind = reductionOp.getKind();

mlir/lib/Dialect/AMX/Transforms/LegalizeForLLVMExport.cpp

 using namespace mlir::amx;
 
 namespace {
 
 /// Maps the 2-dim vector shape to the two 16-bit tile sizes. The first
 /// dimension directly translates into the number of rows of the tiles.
 /// The second dimensions needs to be scaled by the number of bytes.
 std::pair<Value, Value> getTileSizes(ConversionPatternRewriter &rewriter,
-                                     LLVMTypeConverter &typeConverter,
+                                     const LLVMTypeConverter &typeConverter,
                                      VectorType vType, Location loc) {
   Type llvmInt16Type = IntegerType::get(&typeConverter.getContext(), 16);
   unsigned width = vType.getElementType().getIntOrFloatBitWidth();
   assert(llvm::isPowerOf2_64(width) && width >= 8);
   unsigned bytes = width >> 3;
   auto mattr = rewriter.getI16IntegerAttr(vType.getDimSize(0));
   auto nattr = rewriter.getI16IntegerAttr(vType.getDimSize(1) * bytes);
   return std::make_pair(
LogicalResult verifyStride(MemRefType mType) {
   if (failed(getStridesAndOffset(mType, strides, offset)) || strides[last] != 1)
     return failure();
   return success();
 }
 
 /// Maps the 2-dim memref shape to the 64-bit stride. Note that the buffer
 /// shape may "envelop" the actual tile shape, and may be dynamically sized.
 Value getStride(ConversionPatternRewriter &rewriter,
-                LLVMTypeConverter &typeConverter, MemRefType mType, Value base,
-                Location loc) {
+                const LLVMTypeConverter &typeConverter, MemRefType mType,
+                Value base, Location loc) {
   assert(mType.getRank() >= 2);
   int64_t last = mType.getRank() - 1;
   Type llvmInt64Type = IntegerType::get(&typeConverter.getContext(), 64);
   unsigned width = mType.getElementType().getIntOrFloatBitWidth();
   assert(llvm::isPowerOf2_64(width) && width >= 8);
   unsigned bytes = width >> 3;
   if (mType.isDynamicDim(last)) {
     // Dynamic size needs code to compute the stride at runtime.

mlir/lib/Dialect/Arith/Transforms/EmulateUnsupportedFloats.cpp

LogicalResult EmulateFloatPattern::match(Operation *op) const {
   if (op->getNumRegions() != 0)
     return failure();
   return success();
 }
 
 void EmulateFloatPattern::rewrite(Operation *op, ArrayRef<Value> operands,
                                   ConversionPatternRewriter &rewriter) const {
   Location loc = op->getLoc();
-  TypeConverter *converter = getTypeConverter();
+  const TypeConverter *converter = getTypeConverter();
   SmallVector<Type> resultTypes;
   if (failed(converter->convertTypes(op->getResultTypes(), resultTypes))) {
     // Note to anyone looking for this error message: this is a "can't happen".
     // If you're seeing it, there's a bug.
     op->emitOpError("type conversion failed in float emulation");
     return;
   }
   Operation *expandedOp =

mlir/lib/Dialect/SPIRV/Transforms/SPIRVConversion.cpp

 Type SPIRVTypeConverter::getIndexType() const {
   return ::getIndexType(getContext(), options);
 }
 
 MLIRContext *SPIRVTypeConverter::getContext() const {
   return targetEnv.getAttr().getContext();
 }
 
-bool SPIRVTypeConverter::allows(spirv::Capability capability) {
+bool SPIRVTypeConverter::allows(spirv::Capability capability) const {
   return targetEnv.allows(capability);
 }
 
 // TODO: This is a utility function that should probably be exposed by the
 // SPIR-V dialect. Keeping it local till the use case arises.
 static std::optional<int64_t>
 getTypeNumBytes(const SPIRVConversionOptions &options, Type type) {
   if (isa<spirv::ScalarType>(type)) {
Value strideVal = builder.create<spirv::ConstantOp>(
         IntegerAttr::get(integerType, strides[index.index()]));
     Value update = builder.create<spirv::IMulOp>(loc, strideVal, index.value());
     linearizedIndex =
         builder.create<spirv::IAddOp>(loc, linearizedIndex, update);
   }
   return linearizedIndex;
 }
 
-Value mlir::spirv::getVulkanElementPtr(SPIRVTypeConverter &typeConverter,
+Value mlir::spirv::getVulkanElementPtr(const SPIRVTypeConverter &typeConverter,
                                        MemRefType baseType, Value basePtr,
                                        ValueRange indices, Location loc,
                                        OpBuilder &builder) {
   // Get base and offset of the MemRefType and verify they are static.
 
   int64_t offset;
   SmallVector<int64_t, 4> strides;
   if (failed(getStridesAndOffset(baseType, strides, offset)) ||
if (baseType.getRank() == 0) {
     linearizedIndices.push_back(zero);
   } else {
     linearizedIndices.push_back(
         linearizeIndex(indices, strides, offset, indexType, loc, builder));
   }
   return builder.create<spirv::AccessChainOp>(loc, basePtr, linearizedIndices);
 }
 
-Value mlir::spirv::getOpenCLElementPtr(SPIRVTypeConverter &typeConverter,
+Value mlir::spirv::getOpenCLElementPtr(const SPIRVTypeConverter &typeConverter,
                                        MemRefType baseType, Value basePtr,
                                        ValueRange indices, Location loc,
                                        OpBuilder &builder) {
   // Get base and offset of the MemRefType and verify they are static.
 
   int64_t offset;
   SmallVector<int64_t, 4> strides;
   if (failed(getStridesAndOffset(baseType, strides, offset)) ||
if (isa<spirv::ArrayType>(pointeeType)) {
     linearizedIndices.push_back(linearIndex);
     return builder.create<spirv::AccessChainOp>(loc, basePtr,
                                                 linearizedIndices);
   }
   return builder.create<spirv::PtrAccessChainOp>(loc, basePtr, linearIndex,
                                                  linearizedIndices);
 }
 
-Value mlir::spirv::getElementPtr(SPIRVTypeConverter &typeConverter,
+Value mlir::spirv::getElementPtr(const SPIRVTypeConverter &typeConverter,
                                  MemRefType baseType, Value basePtr,
                                  ValueRange indices, Location loc,
                                  OpBuilder &builder) {
 
   if (typeConverter.allows(spirv::Capability::Kernel)) {
     return getOpenCLElementPtr(typeConverter, baseType, basePtr, indices, loc,
                                builder);
   }

mlir/lib/Dialect/X86Vector/Transforms/LegalizeForLLVMExport.cpp

 /// intrinsics based on the bitwidth of their "main" vector element type. This
 /// relies on the to-LLVM-dialect conversion helpers to correctly pack the
 /// results of multi-result intrinsic ops.
 template <typename OpTy, typename Intr32OpTy, typename Intr64OpTy>
 struct LowerToIntrinsic : public OpConversionPattern<OpTy> {
   explicit LowerToIntrinsic(LLVMTypeConverter &converter)
       : OpConversionPattern<OpTy>(converter, &converter.getContext()) {}
 
-  LLVMTypeConverter &getTypeConverter() const {
-    return *static_cast<LLVMTypeConverter *>(
+  const LLVMTypeConverter &getTypeConverter() const {
+    return *static_cast<const LLVMTypeConverter *>(
         OpConversionPattern<OpTy>::getTypeConverter());
   }
 
   LogicalResult
   matchAndRewrite(OpTy op, typename OpTy::Adaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     Type elementType = getSrcVectorElementType<OpTy>(op);
     unsigned bitwidth = elementType.getIntOrFloatBitWidth();

mlir/lib/Transforms/Utils/DialectConversion.cpp

 };
 
 //===----------------------------------------------------------------------===//
 // OpReplacement
 
 /// This class represents one requested operation replacement via 'replaceOp' or
 /// 'eraseOp`.
 struct OpReplacement {
-  OpReplacement(TypeConverter *converter = nullptr) : converter(converter) {}
+  OpReplacement(const TypeConverter *converter = nullptr)
+      : converter(converter) {}
 
   /// An optional type converter that can be used to materialize conversions
   /// between the new and old values if necessary.
-  TypeConverter *converter;
+  const TypeConverter *converter;
 };
 
 //===----------------------------------------------------------------------===//
 // BlockAction
 
 /// The kind of the block action performed during the rewrite.  Actions can be
 /// undone if the conversion fails.
 enum class BlockActionKind {
enum Kind {
     Argument,
 
     /// This materialization materializes a conversion from an illegal type to a
     /// legal one.
     Target
   };
 
   UnresolvedMaterialization(UnrealizedConversionCastOp op = nullptr,
-                            TypeConverter *converter = nullptr,
+                            const TypeConverter *converter = nullptr,
                             Kind kind = Target, Type origOutputType = nullptr)
       : op(op), converterAndKind(converter, kind),
         origOutputType(origOutputType) {}
 
   /// Return the temporary conversion operation inserted for this
   /// materialization.
   UnrealizedConversionCastOp getOp() const { return op; }
 
   /// Return the type converter of this materialization (which may be null).
-  TypeConverter *getConverter() const { return converterAndKind.getPointer(); }
+  const TypeConverter *getConverter() const {
+    return converterAndKind.getPointer();
+  }
 
   /// Return the kind of this materialization.
   Kind getKind() const { return converterAndKind.getInt(); }
 
   /// Set the kind of this materialization.
   void setKind(Kind kind) { converterAndKind.setInt(kind); }
 
   /// Return the original illegal output type of the input values.
   Type getOrigOutputType() const { return origOutputType; }
 
 private:
   /// The unresolved materialization operation created during conversion.
   UnrealizedConversionCastOp op;
 
   /// The corresponding type converter to use when resolving this
   /// materialization, and the kind of this materialization.
-  llvm::PointerIntPair<TypeConverter *, 1, Kind> converterAndKind;
+  llvm::PointerIntPair<const TypeConverter *, 1, Kind> converterAndKind;
 
   /// The original output type. This is only used for argument conversions.
   Type origOutputType;
 };
 } // namespace
 
 /// Build an unresolved materialization operation given an output type and set
 /// of input operands.
 static Value buildUnresolvedMaterialization(
     UnresolvedMaterialization::Kind kind, Block *insertBlock,
     Block::iterator insertPt, Location loc, ValueRange inputs, Type outputType,
-    Type origOutputType, TypeConverter *converter,
+    Type origOutputType, const TypeConverter *converter,
     SmallVectorImpl<UnresolvedMaterialization> &unresolvedMaterializations) {
   // Avoid materializing an unnecessary cast.
   if (inputs.size() == 1 && inputs.front().getType() == outputType)
     return inputs.front();
 
   // Create an unresolved materialization. We use a new OpBuilder to avoid
   // tracking the materialization like we do for other operations.
   OpBuilder builder(insertBlock, insertPt);
   auto convertOp =
       builder.create<UnrealizedConversionCastOp>(loc, outputType, inputs);
   unresolvedMaterializations.emplace_back(convertOp, converter, kind,
                                           origOutputType);
   return convertOp.getResult(0);
 }
 static Value buildUnresolvedArgumentMaterialization(
     PatternRewriter &rewriter, Location loc, ValueRange inputs,
-    Type origOutputType, Type outputType, TypeConverter *converter,
+    Type origOutputType, Type outputType, const TypeConverter *converter,
     SmallVectorImpl<UnresolvedMaterialization> &unresolvedMaterializations) {
   return buildUnresolvedMaterialization(
       UnresolvedMaterialization::Argument, rewriter.getInsertionBlock(),
       rewriter.getInsertionPoint(), loc, inputs, outputType, origOutputType,
       converter, unresolvedMaterializations);
 }
 static Value buildUnresolvedTargetMaterialization(
-    Location loc, Value input, Type outputType, TypeConverter *converter,
+    Location loc, Value input, Type outputType, const TypeConverter *converter,
     SmallVectorImpl<UnresolvedMaterialization> &unresolvedMaterializations) {
   Block *insertBlock = input.getParentBlock();
   Block::iterator insertPt = insertBlock->begin();
   if (OpResult inputRes = dyn_cast<OpResult>(input))
     insertPt = ++inputRes.getOwner()->getIterator();
 
   return buildUnresolvedMaterialization(
       UnresolvedMaterialization::Target, insertBlock, insertPt, loc, input,
struct ConvertedArgInfo {
     /// The cast value that was created to cast from the new arguments to the
     /// old. This only used if 'newArgSize' > 1.
     Value castValue;
   };
 
   /// This structure contains information pertaining to a block that has had its
   /// signature converted.
   struct ConvertedBlockInfo {
-    ConvertedBlockInfo(Block *origBlock, TypeConverter *converter)
+    ConvertedBlockInfo(Block *origBlock, const TypeConverter *converter)
         : origBlock(origBlock), converter(converter) {}
 
     /// The original block that was requested to have its signature converted.
     Block *origBlock;
 
     /// The conversion information for each of the arguments. The information is
     /// std::nullopt if the argument was dropped during conversion.
     SmallVector<std::optional<ConvertedArgInfo>, 1> argInfo;
 
     /// The type converter used to convert the arguments.
-    TypeConverter *converter;
+    const TypeConverter *converter;
   };
 
   /// Return if the signature of the given block has already been converted.
   bool hasBeenConverted(Block *block) const {
     return conversionInfo.count(block) || convertedBlocks.count(block);
   }
 
   /// Set the type converter to use for the given region.
-  void setConverter(Region *region, TypeConverter *typeConverter) {
+  void setConverter(Region *region, const TypeConverter *typeConverter) {
     assert(typeConverter && "expected valid type converter");
     regionToConverter[region] = typeConverter;
   }
 
   /// Return the type converter to use for the given region, or null if there
   /// isn't one.
-  TypeConverter *getConverter(Region *region) {
+  const TypeConverter *getConverter(Region *region) {
     return regionToConverter.lookup(region);
   }
 
   //===--------------------------------------------------------------------===//
   // Rewrite Application
   //===--------------------------------------------------------------------===//
 
   /// Erase any rewrites registered for the blocks within the given operation
struct ArgConverter {
   //===--------------------------------------------------------------------===//
   // Conversion
   //===--------------------------------------------------------------------===//
 
   /// Attempt to convert the signature of the given block, if successful a new
   /// block is returned containing the new arguments. Returns `block` if it did
   /// not require conversion.
   FailureOr<Block *>
-  convertSignature(Block *block, TypeConverter *converter,
+  convertSignature(Block *block, const TypeConverter *converter,
                    ConversionValueMapping &mapping,
                    SmallVectorImpl<BlockArgument> &argReplacements);
 
   /// Apply the given signature conversion on the given block. The new block
   /// containing the updated signature is returned. If no conversions were
   /// necessary, e.g. if the block has no arguments, `block` is returned.
   /// `converter` is used to generate any necessary cast operations that
   /// translate between the origin argument types and those specified in the
   /// signature conversion.
   Block *applySignatureConversion(
-      Block *block, TypeConverter *converter,
+      Block *block, const TypeConverter *converter,
       TypeConverter::SignatureConversion &signatureConversion,
       ConversionValueMapping &mapping,
       SmallVectorImpl<BlockArgument> &argReplacements);
 
   /// Insert a new conversion into the cache.
   void insertConversion(Block *newBlock, ConvertedBlockInfo &&info);
 
   /// A collection of blocks that have had their arguments converted. This is a
   /// map from the new replacement block, back to the original block.
   llvm::MapVector<Block *, ConvertedBlockInfo> conversionInfo;
 
   /// The set of original blocks that were converted.
   DenseSet<Block *> convertedBlocks;
 
   /// A mapping from valid regions, to those containing the original blocks of a
   /// conversion.
   DenseMap<Region *, std::unique_ptr<Region>> regionMapping;
 
   /// A mapping of regions to type converters that should be used when
   /// converting the arguments of blocks within that region.
-  DenseMap<Region *, TypeConverter *> regionToConverter;
+  DenseMap<Region *, const TypeConverter *> regionToConverter;
 
   /// The pattern rewriter to use when materializing conversions.
   PatternRewriter &rewriter;
 
   /// An ordered set of unresolved materializations during conversion.
   SmallVectorImpl<UnresolvedMaterialization> &unresolvedMaterializations;
 };
 } // namespace
LogicalResult ArgConverter::materializeLiveConversions(
   }
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // Conversion
 
 FailureOr<Block *> ArgConverter::convertSignature(
-    Block *block, TypeConverter *converter, ConversionValueMapping &mapping,
+    Block *block, const TypeConverter *converter,
+    ConversionValueMapping &mapping,
     SmallVectorImpl<BlockArgument> &argReplacements) {
   // Check if the block was already converted. If the block is detached,
   // conservatively assume it is going to be deleted.
   if (hasBeenConverted(block) || !block->getParent())
     return block;
   // If a converter wasn't provided, and the block wasn't already converted,
   // there is nothing we can do.
   if (!converter)
     return failure();
 
   // Try to convert the signature for the block with the provided converter.
   if (auto conversion = converter->convertBlockSignature(block))
     return applySignatureConversion(block, converter, *conversion, mapping,
                                     argReplacements);
   return failure();
 }
 
 Block *ArgConverter::applySignatureConversion(
-    Block *block, TypeConverter *converter,
+    Block *block, const TypeConverter *converter,
     TypeConverter::SignatureConversion &signatureConversion,
     ConversionValueMapping &mapping,
     SmallVectorImpl<BlockArgument> &argReplacements) {
   // If no arguments are being changed or added, there is nothing to do.
   unsigned origArgCount = block->getNumArguments();
   auto convertedTypes = signatureConversion.getConvertedTypes();
   if (origArgCount == 0 && convertedTypes.empty())
     return block;
struct ConversionPatternRewriterImpl {
   void markNestedOpsIgnored(Operation *op);
 
   //===--------------------------------------------------------------------===//
   // Type Conversion
   //===--------------------------------------------------------------------===//
 
   /// Convert the signature of the given block.
   FailureOr<Block *> convertBlockSignature(
-      Block *block, TypeConverter *converter,
+      Block *block, const TypeConverter *converter,
       TypeConverter::SignatureConversion *conversion = nullptr);
 
   /// Apply a signature conversion on the given region, using `converter` for
   /// materializations if not null.
   Block *
   applySignatureConversion(Region *region,
                            TypeConverter::SignatureConversion &conversion,
-                           TypeConverter *converter);
+                           const TypeConverter *converter);
 
   /// Convert the types of block arguments within the given region.
   FailureOr<Block *>
-  convertRegionTypes(Region *region, TypeConverter &converter,
+  convertRegionTypes(Region *region, const TypeConverter &converter,
                      TypeConverter::SignatureConversion *entryConversion);
 
   /// Convert the types of non-entry block arguments within the given region.
   LogicalResult convertNonEntryRegionTypes(
-      Region *region, TypeConverter &converter,
+      Region *region, const TypeConverter &converter,
       ArrayRef<TypeConverter::SignatureConversion> blockConversions = {});
 
   //===--------------------------------------------------------------------===//
   // Rewriter Notification Hooks
   //===--------------------------------------------------------------------===//
 
   /// PatternRewriter hook for replacing the results of an operation.
   void notifyOpReplaced(Operation *op, ValueRange newValues);
struct ConversionPatternRewriterImpl {
 
   /// A vector of indices into `replacements` of operations that were replaced
   /// with values with different result types than the original operation, e.g.
   /// 1->N conversion of some kind.
   SmallVector<unsigned, 4> operationsWithChangedResults;
 
   /// The current type converter, or nullptr if no type converter is currently
   /// active.
-  TypeConverter *currentTypeConverter = nullptr;
+  const TypeConverter *currentTypeConverter = nullptr;
 
   /// This allows the user to collect the match failure message.
   function_ref<void(Diagnostic &)> notifyCallback;
 
 #ifndef NDEBUG
   /// A set of operations that have pending updates. This tracking isn't
   /// strictly necessary, and is thus only active during debug builds for extra
   /// verification.
if (llvm::any_of(op->getRegions(),
       ignoredOps.insert(op);
   });
 }
 
 //===----------------------------------------------------------------------===//
 // Type Conversion
 
 FailureOr<Block *> ConversionPatternRewriterImpl::convertBlockSignature(
-    Block *block, TypeConverter *converter,
+    Block *block, const TypeConverter *converter,
     TypeConverter::SignatureConversion *conversion) {
   FailureOr<Block *> result =
       conversion ? argConverter.applySignatureConversion(
                        block, converter, *conversion, mapping, argReplacements)
                  : argConverter.convertSignature(block, converter, mapping,
                                                  argReplacements);
   if (failed(result))
     return failure();
   if (Block *newBlock = *result) {
     if (newBlock != block)
       blockActions.push_back(BlockAction::getTypeConversion(newBlock));
   }
   return result;
 }
 
 Block *ConversionPatternRewriterImpl::applySignatureConversion(
     Region *region, TypeConverter::SignatureConversion &conversion,
-    TypeConverter *converter) {
+    const TypeConverter *converter) {
   if (!region->empty())
     return *convertBlockSignature(&region->front(), converter, &conversion);
   return nullptr;
 }
 
 FailureOr<Block *> ConversionPatternRewriterImpl::convertRegionTypes(
-    Region *region, TypeConverter &converter,
+    Region *region, const TypeConverter &converter,
     TypeConverter::SignatureConversion *entryConversion) {
   argConverter.setConverter(region, &converter);
   if (region->empty())
     return nullptr;
 
   if (failed(convertNonEntryRegionTypes(region, converter)))
     return failure();
 
   FailureOr<Block *> newEntry =
       convertBlockSignature(&region->front(), &converter, entryConversion);
   return newEntry;
 }
 
 LogicalResult ConversionPatternRewriterImpl::convertNonEntryRegionTypes(
-    Region *region, TypeConverter &converter,
+    Region *region, const TypeConverter &converter,
     ArrayRef<TypeConverter::SignatureConversion> blockConversions) {
   argConverter.setConverter(region, &converter);
   if (region->empty())
     return success();
 
   // Convert the arguments of each block within the region.
   int blockIdx = 0;
   assert((blockConversions.empty() ||
void ConversionPatternRewriter::eraseBlock(Block *block) {
   // action and will be actually destroyed when rewrites are applied. This
   // allows us to keep the operations in the block live and undo the removal by
   // re-inserting the block.
   block->getParent()->getBlocks().remove(block);
 }
 
 Block *ConversionPatternRewriter::applySignatureConversion(
     Region *region, TypeConverter::SignatureConversion &conversion,
-    TypeConverter *converter) {
+    const TypeConverter *converter) {
   return impl->applySignatureConversion(region, conversion, converter);
 }
 
 FailureOr<Block *> ConversionPatternRewriter::convertRegionTypes(
-    Region *region, TypeConverter &converter,
+    Region *region, const TypeConverter &converter,
     TypeConverter::SignatureConversion *entryConversion) {
   return impl->convertRegionTypes(region, converter, entryConversion);
 }
 
 LogicalResult ConversionPatternRewriter::convertNonEntryRegionTypes(
-    Region *region, TypeConverter &converter,
+    Region *region, const TypeConverter &converter,
     ArrayRef<TypeConverter::SignatureConversion> blockConversions) {
   return impl->convertNonEntryRegionTypes(region, converter, blockConversions);
 }
 
 void ConversionPatternRewriter::replaceUsesOfBlockArgument(BlockArgument from,
                                                            Value to) {
   LLVM_DEBUG({
     Operation *parentOp = from.getOwner()->getParentOp();
private:
   LogicalResult
   legalizeErasedResult(Operation *op, OpResult result,
                        ConversionPatternRewriterImpl &rewriterImpl);
 
   /// Legalize an operation result that was replaced with a value of a different
   /// type.
   LogicalResult legalizeChangedResultType(
       Operation *op, OpResult result, Value newValue,
-      TypeConverter *replConverter, ConversionPatternRewriter &rewriter,
+      const TypeConverter *replConverter, ConversionPatternRewriter &rewriter,
       ConversionPatternRewriterImpl &rewriterImpl,
       const DenseMap<Value, SmallVector<Value>> &inverseMapping);
 
   /// The legalizer to use when converting operations.
   OperationLegalizer opLegalizer;
 
   /// The conversion mode to use when legalizing operations.
   OpConversionMode mode;
if (remappedValue && remappedValue != opResult) {
                              inverseMapping);
       return success();
     }
   } else {
     // TODO: Avoid materializing other types of conversions here.
   }
 
   // Try to materialize the conversion.
-  if (TypeConverter *converter = mat.getConverter()) {
+  if (const TypeConverter *converter = mat.getConverter()) {
     // FIXME: Determine a suitable insertion location when there are multiple
     // inputs.
     if (inputOperands.size() == 1)
       rewriter.setInsertionPointAfterValue(inputOperands.front());
     else
       rewriter.setInsertionPoint(op);
 
     Value newMaterialization;
while (!worklist.empty()) {
     if (mapIt != inverseMapping.end())
       worklist.append(mapIt->second);
   }
   return nullptr;
 }
 
 LogicalResult OperationConverter::legalizeChangedResultType(
     Operation *op, OpResult result, Value newValue,
-    TypeConverter *replConverter, ConversionPatternRewriter &rewriter,
+    const TypeConverter *replConverter, ConversionPatternRewriter &rewriter,
     ConversionPatternRewriterImpl &rewriterImpl,
     const DenseMap<Value, SmallVector<Value>> &inverseMapping) {
   Operation *liveUser =
       findLiveUserOfReplaced(result, rewriterImpl, inverseMapping);
   if (!liveUser)
     return success();
 
   // Functor used to emit a conversion error for a failed materialization.
TypeConverter::convertTypeAttribute(Type type, Attribute attr) const {
   return std::nullopt;
 }
 
 //===----------------------------------------------------------------------===//
 // FunctionOpInterfaceSignatureConversion
 //===----------------------------------------------------------------------===//
 
 static LogicalResult convertFuncOpTypes(FunctionOpInterface funcOp,
-                                        TypeConverter &typeConverter,
+                                        const TypeConverter &typeConverter,
                                         ConversionPatternRewriter &rewriter) {
   FunctionType type = dyn_cast<FunctionType>(funcOp.getFunctionType());
   if (!type)
     return failure();
 
   // Convert the original function types.
   TypeConverter::SignatureConversion result(type.getNumInputs());
   SmallVector<Type, 1> newResults;

 
 /// Create a default conversion pattern that rewrites the type signature of a
 /// FunctionOpInterface op. This only supports ops which use FunctionType to
 /// represent their type.
 namespace {
 struct FunctionOpInterfaceSignatureConversion : public ConversionPattern {
   FunctionOpInterfaceSignatureConversion(StringRef functionLikeOpName,
                                          MLIRContext *ctx,
-                                         TypeConverter &converter)
+                                         const TypeConverter &converter)
       : ConversionPattern(converter, functionLikeOpName, /*benefit=*/1, ctx) {}
 
   LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> /*operands*/,
                   ConversionPatternRewriter &rewriter) const override {
     FunctionOpInterface funcOp = cast<FunctionOpInterface>(op);
     return convertFuncOpTypes(funcOp, *typeConverter, rewriter);
   }
 };
 
 struct AnyFunctionOpInterfaceSignatureConversion
     : public OpInterfaceConversionPattern<FunctionOpInterface> {
   using OpInterfaceConversionPattern::OpInterfaceConversionPattern;
 
   LogicalResult
   matchAndRewrite(FunctionOpInterface funcOp, ArrayRef<Value> /*operands*/,
                   ConversionPatternRewriter &rewriter) const override {
     return convertFuncOpTypes(funcOp, *typeConverter, rewriter);
   }
 };
 } // namespace
 
 void mlir::populateFunctionOpInterfaceTypeConversionPattern(
     StringRef functionLikeOpName, RewritePatternSet &patterns,
-    TypeConverter &converter) {
+    const TypeConverter &converter) {
   patterns.add<FunctionOpInterfaceSignatureConversion>(
       functionLikeOpName, patterns.getContext(), converter);
 }
 
 void mlir::populateAnyFunctionOpInterfaceTypeConversionPattern(
-    RewritePatternSet &patterns, TypeConverter &converter) {
+    RewritePatternSet &patterns, const TypeConverter &converter) {
   patterns.add<AnyFunctionOpInterfaceSignatureConversion>(
       converter, patterns.getContext());
 }
 
 //===----------------------------------------------------------------------===//
 // ConversionTarget
 //===----------------------------------------------------------------------===//
 
patterns.getPDLPatterns().registerRewriteFunction(
         return pdllConvertValues(
             static_cast<ConversionPatternRewriter &>(rewriter), values);
       });
   patterns.getPDLPatterns().registerRewriteFunction(
       "convertType",
       [](PatternRewriter &rewriter, Type type) -> FailureOr<Type> {
         auto &rewriterImpl =
             static_cast<ConversionPatternRewriter &>(rewriter).getImpl();
-        if (TypeConverter *converter = rewriterImpl.currentTypeConverter) {
+        if (const TypeConverter *converter =
+                rewriterImpl.currentTypeConverter) {
           if (Type newType = converter->convertType(type))
             return newType;
           return failure();
         }
         return type;
       });
   patterns.getPDLPatterns().registerRewriteFunction(
       "convertTypes",
       [](PatternRewriter &rewriter,
          TypeRange types) -> FailureOr<SmallVector<Type>> {
         auto &rewriterImpl =
             static_cast<ConversionPatternRewriter &>(rewriter).getImpl();
-        TypeConverter *converter = rewriterImpl.currentTypeConverter;
+        const TypeConverter *converter = rewriterImpl.currentTypeConverter;
         if (!converter)
           return SmallVector<Type>(types);
 
         SmallVector<Type> remappedTypes;
         if (failed(converter->convertTypes(types, remappedTypes)))
           return failure();
         return std::move(remappedTypes);
       });

mlir/test/lib/Dialect/Test/TestPatterns.cpp

struct TestUndoBlockErase : public ConversionPattern {
   }
 };
 
 //===----------------------------------------------------------------------===//
 // Type-Conversion Rewrite Testing
 
 /// This patterns erases a region operation that has had a type conversion.
 struct TestDropOpSignatureConversion : public ConversionPattern {
-  TestDropOpSignatureConversion(MLIRContext *ctx, TypeConverter &converter)
+  TestDropOpSignatureConversion(MLIRContext *ctx,
+                                const TypeConverter &converter)
       : ConversionPattern(converter, "test.drop_region_op", 1, ctx) {}
   LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     Region &region = op->getRegion(0);
     Block *entry = &region.front();
 
     // Convert the original entry arguments.
-    TypeConverter &converter = *getTypeConverter();
+    const TypeConverter &converter = *getTypeConverter();
     TypeConverter::SignatureConversion result(entry->getNumArguments());
     if (failed(converter.convertSignatureArgs(entry->getArgumentTypes(),
                                               result)) ||
         failed(rewriter.convertRegionTypes(&region, converter, &result)))
       return failure();
 
     // Convert the region signature and just drop the operation.
     rewriter.eraseOp(op);
matchAndRewrite(TestSignatureConversionUndoOp op, OpAdaptor adaptor,
     return failure();
   }
 };
 
 /// Call signature conversion without providing a type converter to handle
 /// materializations.
 struct TestTestSignatureConversionNoConverter
     : public OpConversionPattern<TestSignatureConversionNoConverterOp> {
-  TestTestSignatureConversionNoConverter(TypeConverter &converter,
+  TestTestSignatureConversionNoConverter(const TypeConverter &converter,
                                          MLIRContext *context)
       : OpConversionPattern<TestSignatureConversionNoConverterOp>(context),
         converter(converter) {}
 
   LogicalResult
   matchAndRewrite(TestSignatureConversionNoConverterOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const final {
     Region &region = op->getRegion(0);
     Block *entry = &region.front();
 
     // Convert the original entry arguments.
     TypeConverter::SignatureConversion result(entry->getNumArguments());
     if (failed(
             converter.convertSignatureArgs(entry->getArgumentTypes(), result)))
       return failure();
     rewriter.updateRootInPlace(
         op, [&] { rewriter.applySignatureConversion(&region, result); });
     return success();
   }
 
-  TypeConverter &converter;
+  const TypeConverter &converter;
 };
 
 /// Just forward the operands to the root op. This is essentially a no-op
 /// pattern that is used to trigger target materialization.
 struct TestTypeConsumerForward
     : public OpConversionPattern<TestTypeConsumerOp> {
   using OpConversionPattern<TestTypeConsumerOp>::OpConversionPattern;