[NFC] Use ConvertOpToLLVMPattern instead of ConvertToLLVMPattern.

[jurahul]

Summary

• use ConvertOpToLLVMPattern to avoid explicit casting and in most cases the constructor can be reused to saves a few lines of code.

mlir/include/mlir/Conversion/LinalgToLLVM/LinalgToLLVM.h

 
 namespace mlir {
 class MLIRContext;
 class ModuleOp;
 template <typename T> class OperationPass;
 
 /// Populate the given list with patterns that convert from Linalg to LLVM.
 void populateLinalgToLLVMConversionPatterns(LLVMTypeConverter &converter,
-                                            OwningRewritePatternList &patterns,
-                                            MLIRContext *ctx);
+                                            OwningRewritePatternList &patterns);
 
 /// Create a pass to convert Linalg operations to the LLVMIR dialect.
 std::unique_ptr<OperationPass<ModuleOp>> createConvertLinalgToLLVMPass();
 
 } // namespace mlir
 
 #endif // MLIR_CONVERSION_LINALGTOLLVM_LINALGTOLLVM_H_

mlir/include/mlir/Conversion/OpenMPToLLVM/ConvertOpenMPToLLVM.h

 class LLVMTypeConverter;
 class MLIRContext;
 class ModuleOp;
 template <typename T>
 class OperationPass;
 class OwningRewritePatternList;
 
 /// Populate the given list with patterns that convert from OpenMP to LLVM.
-void populateOpenMPToLLVMConversionPatterns(MLIRContext *context,
-                                            LLVMTypeConverter &converter,
+void populateOpenMPToLLVMConversionPatterns(LLVMTypeConverter &converter,
                                             OwningRewritePatternList &patterns);
 
 /// Create a pass to convert OpenMP operations to the LLVMIR dialect.
 std::unique_ptr<OperationPass<ModuleOp>> createConvertOpenMPToLLVMPass();
 
 } // namespace mlir
 
 #endif // MLIR_CONVERSION_OPENMPTOLLVM_OPENMPTOLLVM_H_

mlir/include/mlir/Conversion/StandardToLLVM/ConvertStandardToLLVM.h

createMemRefDescriptor(Location loc, MemRefType memRefType,
                          ConversionPatternRewriter &rewriter) const;
 };
 
 /// Utility class for operation conversions targeting the LLVM dialect that
 /// match exactly one source operation.
 template <typename SourceOp>
 class ConvertOpToLLVMPattern : public ConvertToLLVMPattern {
 public:
-  ConvertOpToLLVMPattern(LLVMTypeConverter &typeConverter,
+  explicit ConvertOpToLLVMPattern(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 {
     rewrite(cast<SourceOp>(op), operands, rewriter);

mlir/lib/Conversion/AVX512ToLLVM/ConvertAVX512ToLLVM.cpp

 // TODO: Code is currently copy-pasted and adapted from the code
 // 1-1 LLVM conversion. It would better if it were properly exposed in core and
 // reusable.
 /// Basic lowering implementation for one-to-one rewriting from AVX512 Ops to
 /// LLVM Dialect Ops. Convert the type of the result to an LLVM type, pass
 /// operands as is, preserve attributes.
 template <typename SourceOp, typename TargetOp>
 static LogicalResult
-matchAndRewriteOneToOne(const ConvertToLLVMPattern &lowering,
-                        LLVMTypeConverter &typeConverter, Operation *op,
+matchAndRewriteOneToOne(LLVMTypeConverter &typeConverter, Operation *op,
                         ArrayRef<Value> operands,
                         ConversionPatternRewriter &rewriter) {
   unsigned numResults = op->getNumResults();
 
   Type packedType;
   if (numResults != 0) {
     packedType = typeConverter.packFunctionResults(op->getResultTypes());
     if (!packedType)
matchAndRewriteOneToOne(LLVMTypeConverter &typeConverter, Operation *op,
   rewriter.replaceOp(op, results);
   return success();
 }
 
 namespace {
 // TODO: Patterns are too verbose due to the fact that we have 1 op (e.g.
 // MaskRndScaleOp) and different possible target ops. It would be better to take
 // a Functor so that all these conversions become 1-liners.
-struct MaskRndScaleOpPS512Conversion : public ConvertToLLVMPattern {
-  explicit MaskRndScaleOpPS512Conversion(MLIRContext *context,
-                                         LLVMTypeConverter &typeConverter)
-      : ConvertToLLVMPattern(MaskRndScaleOp::getOperationName(), context,
-                             typeConverter) {}
+struct MaskRndScaleOpPS512Conversion
+    : public ConvertOpToLLVMPattern<MaskRndScaleOp> {
+  using ConvertOpToLLVMPattern<MaskRndScaleOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(MaskRndScaleOp op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
-    if (!getSrcVectorElementType(cast<MaskRndScaleOp>(op)).isF32())
+    if (!getSrcVectorElementType(op).isF32())
       return failure();
     return matchAndRewriteOneToOne<MaskRndScaleOp,
                                    LLVM::x86_avx512_mask_rndscale_ps_512>(
-        *this, *getTypeConverter(), op, operands, rewriter);
+        *getTypeConverter(), op, operands, rewriter);
   }
 };
 
-struct MaskRndScaleOpPD512Conversion : public ConvertToLLVMPattern {
-  explicit MaskRndScaleOpPD512Conversion(MLIRContext *context,
-                                         LLVMTypeConverter &typeConverter)
-      : ConvertToLLVMPattern(MaskRndScaleOp::getOperationName(), context,
-                             typeConverter) {}
+struct MaskRndScaleOpPD512Conversion
+    : public ConvertOpToLLVMPattern<MaskRndScaleOp> {
+  using ConvertOpToLLVMPattern<MaskRndScaleOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(MaskRndScaleOp op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
-    if (!getSrcVectorElementType(cast<MaskRndScaleOp>(op)).isF64())
+    if (!getSrcVectorElementType(op).isF64())
       return failure();
     return matchAndRewriteOneToOne<MaskRndScaleOp,
                                    LLVM::x86_avx512_mask_rndscale_pd_512>(
-        *this, *getTypeConverter(), op, operands, rewriter);
+        *getTypeConverter(), op, operands, rewriter);
   }
 };
 
-struct ScaleFOpPS512Conversion : public ConvertToLLVMPattern {
-  explicit ScaleFOpPS512Conversion(MLIRContext *context,
-                                   LLVMTypeConverter &typeConverter)
-      : ConvertToLLVMPattern(MaskScaleFOp::getOperationName(), context,
-                             typeConverter) {}
+struct ScaleFOpPS512Conversion : public ConvertOpToLLVMPattern<MaskScaleFOp> {
+  using ConvertOpToLLVMPattern<MaskScaleFOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(MaskScaleFOp op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
-    if (!getSrcVectorElementType(cast<MaskScaleFOp>(op)).isF32())
+    if (!getSrcVectorElementType(op).isF32())
       return failure();
     return matchAndRewriteOneToOne<MaskScaleFOp,
                                    LLVM::x86_avx512_mask_scalef_ps_512>(
-        *this, *getTypeConverter(), op, operands, rewriter);
+        *getTypeConverter(), op, operands, rewriter);
   }
 };
 
-struct ScaleFOpPD512Conversion : public ConvertToLLVMPattern {
-  explicit ScaleFOpPD512Conversion(MLIRContext *context,
-                                   LLVMTypeConverter &typeConverter)
-      : ConvertToLLVMPattern(MaskScaleFOp::getOperationName(), context,
-                             typeConverter) {}
+struct ScaleFOpPD512Conversion : public ConvertOpToLLVMPattern<MaskScaleFOp> {
+  using ConvertOpToLLVMPattern<MaskScaleFOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(MaskScaleFOp op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
-    if (!getSrcVectorElementType(cast<MaskScaleFOp>(op)).isF64())
+    if (!getSrcVectorElementType(op).isF64())
       return failure();
     return matchAndRewriteOneToOne<MaskScaleFOp,
                                    LLVM::x86_avx512_mask_scalef_pd_512>(
-        *this, *getTypeConverter(), op, operands, rewriter);
+        *getTypeConverter(), op, operands, rewriter);
   }
 };
 } // namespace
 
 /// Populate the given list with patterns that convert from AVX512 to LLVM.
 void mlir::populateAVX512ToLLVMConversionPatterns(
     LLVMTypeConverter &converter, OwningRewritePatternList &patterns) {
-  MLIRContext *ctx = converter.getDialect()->getContext();
   // clang-format off
   patterns.insert<MaskRndScaleOpPS512Conversion,
                   MaskRndScaleOpPD512Conversion,
                   ScaleFOpPS512Conversion,
-                  ScaleFOpPD512Conversion>(ctx, converter);
+                  ScaleFOpPD512Conversion>(converter);
   // clang-format on
 }

mlir/lib/Conversion/GPUCommon/GPUOpsLowering.h

 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
 #include "mlir/Dialect/StandardOps/IR/Ops.h"
 #include "mlir/IR/Builders.h"
 #include "llvm/Support/FormatVariadic.h"
 
 namespace mlir {
 
 template <unsigned AllocaAddrSpace>
-struct GPUFuncOpLowering : ConvertToLLVMPattern {
-  explicit GPUFuncOpLowering(LLVMTypeConverter &typeConverter)
-      : ConvertToLLVMPattern(gpu::GPUFuncOp::getOperationName(),
-                             typeConverter.getDialect()->getContext(),
-                             typeConverter) {}
+struct GPUFuncOpLowering : ConvertOpToLLVMPattern<gpu::GPUFuncOp> {
+  using ConvertOpToLLVMPattern<gpu::GPUFuncOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(gpu::GPUFuncOp gpuFuncOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     assert(operands.empty() && "func op is not expected to have operands");
-    auto gpuFuncOp = cast<gpu::GPUFuncOp>(op);
     Location loc = gpuFuncOp.getLoc();
 
     SmallVector<LLVM::GlobalOp, 3> workgroupBuffers;
     workgroupBuffers.reserve(gpuFuncOp.getNumWorkgroupAttributions());
     for (auto en : llvm::enumerate(gpuFuncOp.getWorkgroupAttributions())) {
       Value attribution = en.value();
 
       auto type = attribution.getType().dyn_cast<MemRefType>();
if (failed(rewriter.convertRegionTypes(
             &llvmFuncOp.getBody(), *typeConverter, &signatureConversion)))
       return failure();
 
     rewriter.eraseOp(gpuFuncOp);
     return success();
   }
 };
 
-struct GPUReturnOpLowering : public ConvertToLLVMPattern {
-  GPUReturnOpLowering(LLVMTypeConverter &typeConverter)
-      : ConvertToLLVMPattern(gpu::ReturnOp::getOperationName(),
-                             typeConverter.getDialect()->getContext(),
-                             typeConverter) {}
+struct GPUReturnOpLowering : public ConvertOpToLLVMPattern<gpu::ReturnOp> {
+  using ConvertOpToLLVMPattern<gpu::ReturnOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(gpu::ReturnOp op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     rewriter.replaceOpWithNewOp<LLVM::ReturnOp>(op, operands);
     return success();
   }
 };
 
 } // namespace mlir
 
 #endif // MLIR_CONVERSION_GPUCOMMON_GPUOPSLOWERING_H_

mlir/lib/Conversion/GPUCommon/IndexIntrinsicsOpLowering.h

 namespace mlir {
 
 // Rewriting that replaces Op with XOp, YOp, or ZOp depending on the dimension
 // that Op operates on.  Op is assumed to return an `std.index` value and
 // XOp, YOp and ZOp are assumed to return an `llvm.i32` value.  Depending on
 // `indexBitwidth`, sign-extend or truncate the resulting value to match the
 // bitwidth expected by the consumers of the value.
 template <typename Op, typename XOp, typename YOp, typename ZOp>
-struct GPUIndexIntrinsicOpLowering : public ConvertToLLVMPattern {
+struct GPUIndexIntrinsicOpLowering : public ConvertOpToLLVMPattern<Op> {
 private:
   enum dimension { X = 0, Y = 1, Z = 2, invalid };
   unsigned indexBitwidth;
 
   static dimension dimensionToIndex(Op op) {
     return StringSwitch<dimension>(op.dimension())
         .Case("x", X)
         .Case("y", Y)
         .Case("z", Z)
         .Default(invalid);
   }
 
 public:
   explicit GPUIndexIntrinsicOpLowering(LLVMTypeConverter &typeConverter)
-      : ConvertToLLVMPattern(Op::getOperationName(),
-                             typeConverter.getDialect()->getContext(),
-                             typeConverter),
+      : ConvertOpToLLVMPattern<Op>(typeConverter),
         indexBitwidth(typeConverter.getIndexTypeBitwidth()) {}
 
   // Convert the kernel arguments to an LLVM type, preserve the rest.
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(Op op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto loc = op->getLoc();
     MLIRContext *context = rewriter.getContext();
     Value newOp;
-    switch (dimensionToIndex(cast<Op>(op))) {
+    switch (dimensionToIndex(op)) {
     case X:
       newOp = rewriter.create<XOp>(loc, LLVM::LLVMType::getInt32Ty(context));
       break;
     case Y:
       newOp = rewriter.create<YOp>(loc, LLVM::LLVMType::getInt32Ty(context));
       break;
     case Z:
       newOp = rewriter.create<ZOp>(loc, LLVM::LLVMType::getInt32Ty(context));

mlir/lib/Conversion/GPUCommon/OpToFuncCallLowering.h

 /// function called and then the result casted back.
 ///
 /// Example with NVVM:
 ///   %exp_f32 = std.exp %arg_f32 : f32
 ///
 /// will be transformed into
 ///   llvm.call @__nv_expf(%arg_f32) : (!llvm.float) -> !llvm.float
 template <typename SourceOp>
-struct OpToFuncCallLowering : public ConvertToLLVMPattern {
+struct OpToFuncCallLowering : public ConvertOpToLLVMPattern<SourceOp> {
 public:
   explicit OpToFuncCallLowering(LLVMTypeConverter &lowering_, StringRef f32Func,
                                 StringRef f64Func)
-      : ConvertToLLVMPattern(SourceOp::getOperationName(),
-                             lowering_.getDialect()->getContext(), lowering_),
-        f32Func(f32Func), f64Func(f64Func) {}
+      : ConvertOpToLLVMPattern<SourceOp>(lowering_), f32Func(f32Func),
+        f64Func(f64Func) {}
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(SourceOp op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     using LLVM::LLVMFuncOp;
     using LLVM::LLVMType;
 
     static_assert(
         std::is_base_of<OpTrait::OneResult<SourceOp>, SourceOp>::value,
         "expected single result op");
 

mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp

 #include "../GPUCommon/IndexIntrinsicsOpLowering.h"
 #include "../GPUCommon/OpToFuncCallLowering.h"
 #include "../PassDetail.h"
 
 using namespace mlir;
 
 namespace {
 
-struct GPUShuffleOpLowering : public ConvertToLLVMPattern {
-  explicit GPUShuffleOpLowering(LLVMTypeConverter &lowering_)
-      : ConvertToLLVMPattern(gpu::ShuffleOp::getOperationName(),
-                             lowering_.getDialect()->getContext(), lowering_) {}
+struct GPUShuffleOpLowering : public ConvertOpToLLVMPattern<gpu::ShuffleOp> {
+  using ConvertOpToLLVMPattern<gpu::ShuffleOp>::ConvertOpToLLVMPattern;
 
   /// Lowers a shuffle to the corresponding NVVM op.
   ///
   /// Convert the `width` argument into an activeMask (a bitmask which specifies
   /// which threads participate in the shuffle) and a maskAndClamp (specifying
   /// the highest lane which participates in the shuffle).
   ///
   ///     %one = llvm.constant(1 : i32) : !llvm.i32
   ///     %shl = llvm.shl %one, %width : !llvm.i32
   ///     %active_mask = llvm.sub %shl, %one : !llvm.i32
   ///     %mask_and_clamp = llvm.sub %width, %one : !llvm.i32
   ///     %shfl = nvvm.shfl.sync.bfly %active_mask, %value, %offset,
   ///         %mask_and_clamp : !llvm<"{ float, i1 }">
   ///     %shfl_value = llvm.extractvalue %shfl[0 : index] :
   ///         !llvm<"{ float, i1 }">
   ///     %shfl_pred = llvm.extractvalue %shfl[1 : index] :
   ///         !llvm<"{ float, i1 }">
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(gpu::ShuffleOp op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     Location loc = op->getLoc();
     gpu::ShuffleOpAdaptor adaptor(operands);
 
     auto valueTy = adaptor.value().getType().cast<LLVM::LLVMType>();
     auto int32Type = LLVM::LLVMType::getInt32Ty(rewriter.getContext());
     auto predTy = LLVM::LLVMType::getInt1Ty(rewriter.getContext());
     auto resultTy =

mlir/lib/Conversion/LinalgToLLVM/LinalgToLLVM.cpp

 private:
   OpBuilder &rewriter() { return ScopedContext::getBuilderRef(); }
   Location loc() { return ScopedContext::getLocation(); }
 
   MemRefDescriptor d;
 };
 
 // RangeOp creates a new range descriptor.
-class RangeOpConversion : public ConvertToLLVMPattern {
+class RangeOpConversion : public ConvertOpToLLVMPattern<RangeOp> {
 public:
-  explicit RangeOpConversion(MLIRContext *context, LLVMTypeConverter &lowering_)
-      : ConvertToLLVMPattern(RangeOp::getOperationName(), context, lowering_) {}
+  using ConvertOpToLLVMPattern<RangeOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(RangeOp rangeOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
-    auto rangeOp = cast<RangeOp>(op);
     auto rangeDescriptorTy = convertRangeType(
         rangeOp.getType().cast<RangeType>(), *getTypeConverter());
 
-    edsc::ScopedContext context(rewriter, op->getLoc());
+    edsc::ScopedContext context(rewriter, rangeOp->getLoc());
 
     // Fill in an aggregate value of the descriptor.
     RangeOpAdaptor adaptor(operands);
     Value desc = llvm_undef(rangeDescriptorTy);
     desc = llvm_insertvalue(desc, adaptor.min(), rewriter.getI64ArrayAttr(0));
     desc = llvm_insertvalue(desc, adaptor.max(), rewriter.getI64ArrayAttr(1));
     desc = llvm_insertvalue(desc, adaptor.step(), rewriter.getI64ArrayAttr(2));
-    rewriter.replaceOp(op, desc);
+    rewriter.replaceOp(rangeOp, desc);
     return success();
   }
 };
 
 // ReshapeOp creates a new view descriptor of the proper rank.
 // For now, the only conversion supported is for target MemRef with static sizes
 // and strides.
-class ReshapeOpConversion : public ConvertToLLVMPattern {
+class ReshapeOpConversion : public ConvertOpToLLVMPattern<ReshapeOp> {
 public:
-  explicit ReshapeOpConversion(MLIRContext *context,
-                               LLVMTypeConverter &lowering_)
-      : ConvertToLLVMPattern(ReshapeOp::getOperationName(), context,
-                             lowering_) {}
+  using ConvertOpToLLVMPattern<ReshapeOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(ReshapeOp reshapeOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
-    auto reshapeOp = cast<ReshapeOp>(op);
     MemRefType dstType = reshapeOp.getResultType();
 
     if (!dstType.hasStaticShape())
       return failure();
 
     int64_t offset;
     SmallVector<int64_t, 4> strides;
     auto res = getStridesAndOffset(dstType, strides, offset);
     if (failed(res) || llvm::any_of(strides, [](int64_t val) {
           return ShapedType::isDynamicStrideOrOffset(val);
         }))
       return failure();
 
-    edsc::ScopedContext context(rewriter, op->getLoc());
+    edsc::ScopedContext context(rewriter, reshapeOp->getLoc());
     ReshapeOpAdaptor adaptor(operands);
     BaseViewConversionHelper baseDesc(adaptor.src());
     BaseViewConversionHelper desc(typeConverter->convertType(dstType));
     desc.setAllocatedPtr(baseDesc.allocatedPtr());
     desc.setAlignedPtr(baseDesc.alignedPtr());
     desc.setOffset(baseDesc.offset());
     for (auto en : llvm::enumerate(dstType.getShape()))
       desc.setConstantSize(en.index(), en.value());
     for (auto en : llvm::enumerate(strides))
       desc.setConstantStride(en.index(), en.value());
-    rewriter.replaceOp(op, {desc});
+    rewriter.replaceOp(reshapeOp, {desc});
     return success();
   }
 };
 
 /// Conversion pattern that transforms a linalg.slice op into:
 ///   1. An "undef" value for the ViewDescriptor.
 ///   2. Updates to the ViewDescriptor to introduce the data ptr, offset, size
 ///      and stride corresponding to the region of memory within the bounds of
 ///      the parent view.
 /// The linalg.slice op is replaced by the alloca'ed pointer.
-class SliceOpConversion : public ConvertToLLVMPattern {
+class SliceOpConversion : public ConvertOpToLLVMPattern<SliceOp> {
 public:
-  explicit SliceOpConversion(MLIRContext *context, LLVMTypeConverter &lowering_)
-      : ConvertToLLVMPattern(SliceOp::getOperationName(), context, lowering_) {}
+  using ConvertOpToLLVMPattern<SliceOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(SliceOp sliceOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
-    edsc::ScopedContext context(rewriter, op->getLoc());
+    edsc::ScopedContext context(rewriter, sliceOp->getLoc());
     SliceOpAdaptor adaptor(operands);
     BaseViewConversionHelper baseDesc(adaptor.view());
 
-    auto sliceOp = cast<SliceOp>(op);
     auto memRefType = sliceOp.getBaseViewType();
     auto int64Ty = typeConverter->convertType(rewriter.getIntegerType(64))
                        .cast<LLVM::LLVMType>();
 
     BaseViewConversionHelper desc(
         typeConverter->convertType(sliceOp.getShapedType()));
 
     // TODO: extract sizes and emit asserts.
matchAndRewrite(SliceOp sliceOp, ArrayRef<Value> operands,
     desc.setAllocatedPtr(baseDesc.allocatedPtr());
     desc.setAlignedPtr(baseDesc.alignedPtr());
 
     // Insert base offset.
     desc.setOffset(baseOffset);
 
     // Corner case, no sizes or strides: early return the descriptor.
     if (sliceOp.getShapedType().getRank() == 0)
-      return rewriter.replaceOp(op, {desc}), success();
+      return rewriter.replaceOp(sliceOp, {desc}), success();
 
     Value zero = llvm_constant(
         int64Ty, rewriter.getIntegerAttr(rewriter.getIndexType(), 0));
     // Compute and insert view sizes (max - min along the range) and strides.
     // Skip the non-range operands as they will be projected away from the view.
     int numNewDims = 0;
     for (auto en : llvm::enumerate(sliceOp.indexings())) {
       Value indexing = en.value();
for (auto en : llvm::enumerate(sliceOp.indexings())) {
             llvm_select(llvm_icmp(ICmpPredicate::slt, size, zero), zero, size);
         Value stride = llvm_mul(strides[rank], step);
         desc.setSize(numNewDims, size);
         desc.setStride(numNewDims, stride);
         ++numNewDims;
       }
     }
 
-    rewriter.replaceOp(op, {desc});
+    rewriter.replaceOp(sliceOp, {desc});
     return success();
   }
 };
 
 // YieldOp produces and LLVM::ReturnOp.
-class YieldOpConversion : public ConvertToLLVMPattern {
+class YieldOpConversion : public ConvertOpToLLVMPattern<linalg::YieldOp> {
 public:
-  explicit YieldOpConversion(MLIRContext *context, LLVMTypeConverter &lowering_)
-      : ConvertToLLVMPattern(linalg::YieldOp::getOperationName(), context,
-                             lowering_) {}
+  using ConvertOpToLLVMPattern<linalg::YieldOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(linalg::YieldOp op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     rewriter.replaceOpWithNewOp<LLVM::ReturnOp>(op, operands);
     return success();
   }
 };
 } // namespace
 
 /// Populate the given list with patterns that convert from Linalg to LLVM.
 void mlir::populateLinalgToLLVMConversionPatterns(
-    LLVMTypeConverter &converter, OwningRewritePatternList &patterns,
-    MLIRContext *ctx) {
+    LLVMTypeConverter &converter, OwningRewritePatternList &patterns) {
   patterns.insert<RangeOpConversion, ReshapeOpConversion, SliceOpConversion,
-                  YieldOpConversion>(ctx, converter);
+                  YieldOpConversion>(converter);
 
   // Populate the type conversions for the linalg types.
   converter.addConversion(
       [&](RangeType type) { return convertRangeType(type, converter); });
 }
 
 namespace {
 struct ConvertLinalgToLLVMPass
void ConvertLinalgToLLVMPass::runOnOperation() {
   OwningRewritePatternList patterns;
   LLVMTypeConverter converter(&getContext());
   populateAffineToStdConversionPatterns(patterns, &getContext());
   populateLoopToStdConversionPatterns(patterns, &getContext());
   populateStdToLLVMConversionPatterns(converter, patterns);
   populateVectorToSCFConversionPatterns(patterns, &getContext());
   populateVectorToLLVMMatrixConversionPatterns(converter, patterns);
   populateVectorToLLVMConversionPatterns(converter, patterns);
-  populateLinalgToLLVMConversionPatterns(converter, patterns, &getContext());
+  populateLinalgToLLVMConversionPatterns(converter, patterns);
 
   LLVMConversionTarget target(getContext());
   target.addLegalOp<ModuleOp, ModuleTerminatorOp>();
   if (failed(applyFullConversion(module, target, std::move(patterns))))
     signalPassFailure();
 }
 
 std::unique_ptr<OperationPass<ModuleOp>> mlir::createConvertLinalgToLLVMPass() {
   return std::make_unique<ConvertLinalgToLLVMPass>();
 }

mlir/lib/Conversion/OpenMPToLLVM/OpenMPToLLVM.cpp

 using namespace mlir;
 
 namespace {
 /// A pattern that converts the region arguments in a single-region OpenMP
 /// operation to the LLVM dialect. The body of the region is not modified and is
 /// expected to either be processed by the conversion infrastructure or already
 /// contain ops compatible with LLVM dialect types.
 template <typename OpType>
-struct RegionOpConversion : public ConvertToLLVMPattern {
-  explicit RegionOpConversion(MLIRContext *context,
-                              LLVMTypeConverter &typeConverter)
-      : ConvertToLLVMPattern(OpType::getOperationName(), context,
-                             typeConverter) {}
+struct RegionOpConversion : public ConvertOpToLLVMPattern<OpType> {
+  using ConvertOpToLLVMPattern<OpType>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(OpType curOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
-    auto curOp = cast<OpType>(op);
     auto newOp = rewriter.create<OpType>(curOp.getLoc(), TypeRange(), operands,
                                          curOp.getAttrs());
     rewriter.inlineRegionBefore(curOp.region(), newOp.region(),
                                 newOp.region().end());
-    if (failed(rewriter.convertRegionTypes(&newOp.region(), *typeConverter)))
+    if (failed(rewriter.convertRegionTypes(&newOp.region(),
+                                           *this->getTypeConverter())))
       return failure();
 
-    rewriter.eraseOp(op);
+    rewriter.eraseOp(curOp);
     return success();
   }
 };
 } // namespace
 
 void mlir::populateOpenMPToLLVMConversionPatterns(
-    MLIRContext *context, LLVMTypeConverter &converter,
-    OwningRewritePatternList &patterns) {
+    LLVMTypeConverter &converter, OwningRewritePatternList &patterns) {
   patterns.insert<RegionOpConversion<omp::ParallelOp>,
-                  RegionOpConversion<omp::WsLoopOp>>(context, converter);
+                  RegionOpConversion<omp::WsLoopOp>>(converter);
 }
 
 namespace {
 struct ConvertOpenMPToLLVMPass
     : public ConvertOpenMPToLLVMBase<ConvertOpenMPToLLVMPass> {
   void runOnOperation() override;
 };
 } // namespace
 
 void ConvertOpenMPToLLVMPass::runOnOperation() {
   auto module = getOperation();
-  MLIRContext *context = &getContext();
 
   // Convert to OpenMP operations with LLVM IR dialect
   OwningRewritePatternList patterns;
   LLVMTypeConverter converter(&getContext());
   populateStdToLLVMConversionPatterns(converter, patterns);
-  populateOpenMPToLLVMConversionPatterns(context, converter, patterns);
+  populateOpenMPToLLVMConversionPatterns(converter, patterns);
 
   LLVMConversionTarget target(getContext());
   target.addDynamicallyLegalOp<omp::ParallelOp, omp::WsLoopOp>(
       [&](Operation *op) { return converter.isLegal(&op->getRegion(0)); });
   target.addLegalOp<omp::TerminatorOp, omp::TaskyieldOp, omp::FlushOp,
                     omp::BarrierOp, omp::TaskwaitOp>();
   if (failed(applyPartialConversion(module, target, std::move(patterns))))
     signalPassFailure();
 }
 
 std::unique_ptr<OperationPass<ModuleOp>> mlir::createConvertOpenMPToLLVMPass() {
   return std::make_unique<ConvertOpenMPToLLVMPass>();
 }

mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVM.cpp

static TransferWriteOpAdaptor getTransferOpAdapter(TransferWriteOp xferOp,
                                                    ArrayRef<Value> operands) {
   return TransferWriteOpAdaptor(operands);
 }
 
 namespace {
 
 /// Conversion pattern for a vector.matrix_multiply.
 /// This is lowered directly to the proper llvm.intr.matrix.multiply.
-class VectorMatmulOpConversion : public ConvertToLLVMPattern {
+class VectorMatmulOpConversion
+    : public ConvertOpToLLVMPattern<vector::MatmulOp> {
 public:
-  explicit VectorMatmulOpConversion(MLIRContext *context,
-                                    LLVMTypeConverter &typeConverter)
-      : ConvertToLLVMPattern(vector::MatmulOp::getOperationName(), context,
-                             typeConverter) {}
+  using ConvertOpToLLVMPattern<vector::MatmulOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(vector::MatmulOp matmulOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
-    auto matmulOp = cast<vector::MatmulOp>(op);
     auto adaptor = vector::MatmulOpAdaptor(operands);
     rewriter.replaceOpWithNewOp<LLVM::MatrixMultiplyOp>(
-        op, typeConverter->convertType(matmulOp.res().getType()), adaptor.lhs(),
-        adaptor.rhs(), matmulOp.lhs_rows(), matmulOp.lhs_columns(),
-        matmulOp.rhs_columns());
+        matmulOp, typeConverter->convertType(matmulOp.res().getType()),
+        adaptor.lhs(), adaptor.rhs(), matmulOp.lhs_rows(),
+        matmulOp.lhs_columns(), matmulOp.rhs_columns());
     return success();
   }
 };
 
 /// Conversion pattern for a vector.flat_transpose.
 /// This is lowered directly to the proper llvm.intr.matrix.transpose.
-class VectorFlatTransposeOpConversion : public ConvertToLLVMPattern {
+class VectorFlatTransposeOpConversion
+    : public ConvertOpToLLVMPattern<vector::FlatTransposeOp> {
 public:
-  explicit VectorFlatTransposeOpConversion(MLIRContext *context,
-                                           LLVMTypeConverter &typeConverter)
-      : ConvertToLLVMPattern(vector::FlatTransposeOp::getOperationName(),
-                             context, typeConverter) {}
+  using ConvertOpToLLVMPattern<vector::FlatTransposeOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(vector::FlatTransposeOp transOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
-    auto transOp = cast<vector::FlatTransposeOp>(op);
     auto adaptor = vector::FlatTransposeOpAdaptor(operands);
     rewriter.replaceOpWithNewOp<LLVM::MatrixTransposeOp>(
         transOp, typeConverter->convertType(transOp.res().getType()),
         adaptor.matrix(), transOp.rows(), transOp.columns());
     return success();
   }
 };
 
 /// Conversion pattern for a vector.maskedload.
-class VectorMaskedLoadOpConversion : public ConvertToLLVMPattern {
+class VectorMaskedLoadOpConversion
+    : public ConvertOpToLLVMPattern<vector::MaskedLoadOp> {
 public:
-  explicit VectorMaskedLoadOpConversion(MLIRContext *context,
-                                        LLVMTypeConverter &typeConverter)
-      : ConvertToLLVMPattern(vector::MaskedLoadOp::getOperationName(), context,
-                             typeConverter) {}
+  using ConvertOpToLLVMPattern<vector::MaskedLoadOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(vector::MaskedLoadOp load, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
-    auto loc = op->getLoc();
-    auto load = cast<vector::MaskedLoadOp>(op);
+    auto loc = load->getLoc();
     auto adaptor = vector::MaskedLoadOpAdaptor(operands);
 
     // Resolve alignment.
     unsigned align;
     if (failed(getMemRefAlignment(*getTypeConverter(), load, align)))
       return failure();
 
     auto vtype = typeConverter->convertType(load.getResultVectorType());
     Value ptr;
     if (failed(getBasePtr(rewriter, loc, adaptor.base(), load.getMemRefType(),
                           vtype, ptr)))
       return failure();
 
     rewriter.replaceOpWithNewOp<LLVM::MaskedLoadOp>(
         load, vtype, ptr, adaptor.mask(), adaptor.pass_thru(),
         rewriter.getI32IntegerAttr(align));
     return success();
   }
 };
 
 /// Conversion pattern for a vector.maskedstore.
-class VectorMaskedStoreOpConversion : public ConvertToLLVMPattern {
+class VectorMaskedStoreOpConversion
+    : public ConvertOpToLLVMPattern<vector::MaskedStoreOp> {
 public:
-  explicit VectorMaskedStoreOpConversion(MLIRContext *context,
-                                         LLVMTypeConverter &typeConverter)
-      : ConvertToLLVMPattern(vector::MaskedStoreOp::getOperationName(), context,
-                             typeConverter) {}
+  using ConvertOpToLLVMPattern<vector::MaskedStoreOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(vector::MaskedStoreOp store, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
-    auto loc = op->getLoc();
-    auto store = cast<vector::MaskedStoreOp>(op);
+    auto loc = store->getLoc();
     auto adaptor = vector::MaskedStoreOpAdaptor(operands);
 
     // Resolve alignment.
     unsigned align;
     if (failed(getMemRefAlignment(*getTypeConverter(), store, align)))
       return failure();
 
     auto vtype = typeConverter->convertType(store.getValueVectorType());
     Value ptr;
     if (failed(getBasePtr(rewriter, loc, adaptor.base(), store.getMemRefType(),
                           vtype, ptr)))
       return failure();
 
     rewriter.replaceOpWithNewOp<LLVM::MaskedStoreOp>(
         store, adaptor.value(), ptr, adaptor.mask(),
         rewriter.getI32IntegerAttr(align));
     return success();
   }
 };
 
 /// Conversion pattern for a vector.gather.
-class VectorGatherOpConversion : public ConvertToLLVMPattern {
+class VectorGatherOpConversion
+    : public ConvertOpToLLVMPattern<vector::GatherOp> {
 public:
-  explicit VectorGatherOpConversion(MLIRContext *context,
-                                    LLVMTypeConverter &typeConverter)
-      : ConvertToLLVMPattern(vector::GatherOp::getOperationName(), context,
-                             typeConverter) {}
+  using ConvertOpToLLVMPattern<vector::GatherOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(vector::GatherOp gather, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
-    auto loc = op->getLoc();
-    auto gather = cast<vector::GatherOp>(op);
+    auto loc = gather->getLoc();
     auto adaptor = vector::GatherOpAdaptor(operands);
 
     // Resolve alignment.
     unsigned align;
     if (failed(getMemRefAlignment(*getTypeConverter(), gather, align)))
       return failure();
 
     // Get index ptrs.
     VectorType vType = gather.getResultVectorType();
     Type iType = gather.getIndicesVectorType().getElementType();
     Value ptrs;
     if (failed(getIndexedPtrs(rewriter, loc, adaptor.base(), adaptor.indices(),
                               gather.getMemRefType(), vType, iType, ptrs)))
       return failure();
 
     // Replace with the gather intrinsic.
     rewriter.replaceOpWithNewOp<LLVM::masked_gather>(
         gather, typeConverter->convertType(vType), ptrs, adaptor.mask(),
         adaptor.pass_thru(), rewriter.getI32IntegerAttr(align));
     return success();
   }
 };
 
 /// Conversion pattern for a vector.scatter.
-class VectorScatterOpConversion : public ConvertToLLVMPattern {
+class VectorScatterOpConversion
+    : public ConvertOpToLLVMPattern<vector::ScatterOp> {
 public:
-  explicit VectorScatterOpConversion(MLIRContext *context,
-                                     LLVMTypeConverter &typeConverter)
-      : ConvertToLLVMPattern(vector::ScatterOp::getOperationName(), context,
-                             typeConverter) {}
+  using ConvertOpToLLVMPattern<vector::ScatterOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(vector::ScatterOp scatter, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
-    auto loc = op->getLoc();
-    auto scatter = cast<vector::ScatterOp>(op);
+    auto loc = scatter->getLoc();
     auto adaptor = vector::ScatterOpAdaptor(operands);
 
     // Resolve alignment.
     unsigned align;
     if (failed(getMemRefAlignment(*getTypeConverter(), scatter, align)))
       return failure();
 
     // Get index ptrs.
     VectorType vType = scatter.getValueVectorType();
     Type iType = scatter.getIndicesVectorType().getElementType();
     Value ptrs;
     if (failed(getIndexedPtrs(rewriter, loc, adaptor.base(), adaptor.indices(),
                               scatter.getMemRefType(), vType, iType, ptrs)))
       return failure();
 
     // Replace with the scatter intrinsic.
     rewriter.replaceOpWithNewOp<LLVM::masked_scatter>(
         scatter, adaptor.value(), ptrs, adaptor.mask(),
         rewriter.getI32IntegerAttr(align));
     return success();
   }
 };
 
 /// Conversion pattern for a vector.expandload.
-class VectorExpandLoadOpConversion : public ConvertToLLVMPattern {
+class VectorExpandLoadOpConversion
+    : public ConvertOpToLLVMPattern<vector::ExpandLoadOp> {
 public:
-  explicit VectorExpandLoadOpConversion(MLIRContext *context,
-                                        LLVMTypeConverter &typeConverter)
-      : ConvertToLLVMPattern(vector::ExpandLoadOp::getOperationName(), context,
-                             typeConverter) {}
+  using ConvertOpToLLVMPattern<vector::ExpandLoadOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(vector::ExpandLoadOp expand, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
-    auto loc = op->getLoc();
-    auto expand = cast<vector::ExpandLoadOp>(op);
+    auto loc = expand->getLoc();
     auto adaptor = vector::ExpandLoadOpAdaptor(operands);
 
     Value ptr;
     if (failed(getBasePtr(rewriter, loc, adaptor.base(), expand.getMemRefType(),
                           ptr)))
       return failure();
 
     auto vType = expand.getResultVectorType();
     rewriter.replaceOpWithNewOp<LLVM::masked_expandload>(
-        op, typeConverter->convertType(vType), ptr, adaptor.mask(),
+        expand, typeConverter->convertType(vType), ptr, adaptor.mask(),
         adaptor.pass_thru());
     return success();
   }
 };
 
 /// Conversion pattern for a vector.compressstore.
-class VectorCompressStoreOpConversion : public ConvertToLLVMPattern {
+class VectorCompressStoreOpConversion
+    : public ConvertOpToLLVMPattern<vector::CompressStoreOp> {
 public:
-  explicit VectorCompressStoreOpConversion(MLIRContext *context,
-                                           LLVMTypeConverter &typeConverter)
-      : ConvertToLLVMPattern(vector::CompressStoreOp::getOperationName(),
-                             context, typeConverter) {}
+  using ConvertOpToLLVMPattern<vector::CompressStoreOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(vector::CompressStoreOp compress, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
-    auto loc = op->getLoc();
-    auto compress = cast<vector::CompressStoreOp>(op);
+    auto loc = compress->getLoc();
     auto adaptor = vector::CompressStoreOpAdaptor(operands);
 
     Value ptr;
     if (failed(getBasePtr(rewriter, loc, adaptor.base(),
                           compress.getMemRefType(), ptr)))
       return failure();
 
     rewriter.replaceOpWithNewOp<LLVM::masked_compressstore>(
-        op, adaptor.value(), ptr, adaptor.mask());
+        compress, adaptor.value(), ptr, adaptor.mask());
     return success();
   }
 };
 
 /// Conversion pattern for all vector reductions.
-class VectorReductionOpConversion : public ConvertToLLVMPattern {
+class VectorReductionOpConversion
+    : public ConvertOpToLLVMPattern<vector::ReductionOp> {
 public:
-  explicit VectorReductionOpConversion(MLIRContext *context,
-                                       LLVMTypeConverter &typeConverter,
+  explicit VectorReductionOpConversion(LLVMTypeConverter &typeConv,
                                        bool reassociateFPRed)
-      : ConvertToLLVMPattern(vector::ReductionOp::getOperationName(), context,
-                             typeConverter),
+      : ConvertOpToLLVMPattern<vector::ReductionOp>(typeConv),
         reassociateFPReductions(reassociateFPRed) {}
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(vector::ReductionOp reductionOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
-    auto reductionOp = cast<vector::ReductionOp>(op);
     auto kind = reductionOp.kind();
     Type eltType = reductionOp.dest().getType();
     Type llvmType = typeConverter->convertType(eltType);
     if (eltType.isIntOrIndex()) {
       // Integer reductions: add/mul/min/max/and/or/xor.
       if (kind == "add")
         rewriter.replaceOpWithNewOp<LLVM::vector_reduce_add>(
-            op, llvmType, operands[0]);
+            reductionOp, llvmType, operands[0]);
       else if (kind == "mul")
         rewriter.replaceOpWithNewOp<LLVM::vector_reduce_mul>(
-            op, llvmType, operands[0]);
+            reductionOp, llvmType, operands[0]);
       else if (kind == "min" &&
                (eltType.isIndex() || eltType.isUnsignedInteger()))
         rewriter.replaceOpWithNewOp<LLVM::vector_reduce_umin>(
-            op, llvmType, operands[0]);
+            reductionOp, llvmType, operands[0]);
       else if (kind == "min")
         rewriter.replaceOpWithNewOp<LLVM::vector_reduce_smin>(
-            op, llvmType, operands[0]);
+            reductionOp, llvmType, operands[0]);
       else if (kind == "max" &&
                (eltType.isIndex() || eltType.isUnsignedInteger()))
         rewriter.replaceOpWithNewOp<LLVM::vector_reduce_umax>(
-            op, llvmType, operands[0]);
+            reductionOp, llvmType, operands[0]);
       else if (kind == "max")
         rewriter.replaceOpWithNewOp<LLVM::vector_reduce_smax>(
-            op, llvmType, operands[0]);
+            reductionOp, llvmType, operands[0]);
       else if (kind == "and")
         rewriter.replaceOpWithNewOp<LLVM::vector_reduce_and>(
-            op, llvmType, operands[0]);
+            reductionOp, llvmType, operands[0]);
       else if (kind == "or")
         rewriter.replaceOpWithNewOp<LLVM::vector_reduce_or>(
-            op, llvmType, operands[0]);
+            reductionOp, llvmType, operands[0]);
       else if (kind == "xor")
         rewriter.replaceOpWithNewOp<LLVM::vector_reduce_xor>(
-            op, llvmType, operands[0]);
+            reductionOp, llvmType, operands[0]);
       else
         return failure();
       return success();
     }
 
     if (!eltType.isa<FloatType>())
       return failure();
 
     // Floating-point reductions: add/mul/min/max
     if (kind == "add") {
       // Optional accumulator (or zero).
       Value acc = operands.size() > 1 ? operands[1]
                                       : rewriter.create<LLVM::ConstantOp>(
-                                            op->getLoc(), llvmType,
+                                            reductionOp->getLoc(), llvmType,
                                             rewriter.getZeroAttr(eltType));
       rewriter.replaceOpWithNewOp<LLVM::vector_reduce_fadd>(
-          op, llvmType, acc, operands[0],
+          reductionOp, llvmType, acc, operands[0],
           rewriter.getBoolAttr(reassociateFPReductions));
     } else if (kind == "mul") {
       // Optional accumulator (or one).
       Value acc = operands.size() > 1
                       ? operands[1]
                       : rewriter.create<LLVM::ConstantOp>(
-                            op->getLoc(), llvmType,
+                            reductionOp->getLoc(), llvmType,
                             rewriter.getFloatAttr(eltType, 1.0));
       rewriter.replaceOpWithNewOp<LLVM::vector_reduce_fmul>(
-          op, llvmType, acc, operands[0],
+          reductionOp, llvmType, acc, operands[0],
           rewriter.getBoolAttr(reassociateFPReductions));
     } else if (kind == "min")
-      rewriter.replaceOpWithNewOp<LLVM::vector_reduce_fmin>(op, llvmType,
-                                                            operands[0]);
+      rewriter.replaceOpWithNewOp<LLVM::vector_reduce_fmin>(
+          reductionOp, llvmType, operands[0]);
     else if (kind == "max")
-      rewriter.replaceOpWithNewOp<LLVM::vector_reduce_fmax>(op, llvmType,
-                                                            operands[0]);
+      rewriter.replaceOpWithNewOp<LLVM::vector_reduce_fmax>(
+          reductionOp, llvmType, operands[0]);
     else
       return failure();
     return success();
   }
 
 private:
   const bool reassociateFPReductions;
 };
 
 /// Conversion pattern for a vector.create_mask (1-D only).
-class VectorCreateMaskOpConversion : public ConvertToLLVMPattern {
+class VectorCreateMaskOpConversion
+    : public ConvertOpToLLVMPattern<vector::CreateMaskOp> {
 public:
-  explicit VectorCreateMaskOpConversion(MLIRContext *context,
-                                        LLVMTypeConverter &typeConverter,
+  explicit VectorCreateMaskOpConversion(LLVMTypeConverter &typeConv,
                                         bool enableIndexOpt)
-      : ConvertToLLVMPattern(vector::CreateMaskOp::getOperationName(), context,
-                             typeConverter),
+      : ConvertOpToLLVMPattern<vector::CreateMaskOp>(typeConv),
         enableIndexOptimizations(enableIndexOpt) {}
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(vector::CreateMaskOp op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto dstType = op->getResult(0).getType().cast<VectorType>();
     int64_t rank = dstType.getRank();
     if (rank == 1) {
       rewriter.replaceOp(
           op, buildVectorComparison(rewriter, op, enableIndexOptimizations,
                                     dstType.getDimSize(0), operands[0]));
       return success();
     }
     return failure();
   }
 
 private:
   const bool enableIndexOptimizations;
 };
 
-class VectorShuffleOpConversion : public ConvertToLLVMPattern {
+class VectorShuffleOpConversion
+    : public ConvertOpToLLVMPattern<vector::ShuffleOp> {
 public:
-  explicit VectorShuffleOpConversion(MLIRContext *context,
-                                     LLVMTypeConverter &typeConverter)
-      : ConvertToLLVMPattern(vector::ShuffleOp::getOperationName(), context,
-                             typeConverter) {}
+  using ConvertOpToLLVMPattern<vector::ShuffleOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(vector::ShuffleOp shuffleOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
-    auto loc = op->getLoc();
+    auto loc = shuffleOp->getLoc();
     auto adaptor = vector::ShuffleOpAdaptor(operands);
-    auto shuffleOp = cast<vector::ShuffleOp>(op);
     auto v1Type = shuffleOp.getV1VectorType();
     auto v2Type = shuffleOp.getV2VectorType();
     auto vectorType = shuffleOp.getVectorType();
     Type llvmType = typeConverter->convertType(vectorType);
     auto maskArrayAttr = shuffleOp.mask();
 
     // Bail if result type cannot be lowered.
     if (!llvmType)
       return failure();
 
     // Get rank and dimension sizes.
     int64_t rank = vectorType.getRank();
     assert(v1Type.getRank() == rank);
     assert(v2Type.getRank() == rank);
     int64_t v1Dim = v1Type.getDimSize(0);
 
     // For rank 1, where both operands have *exactly* the same vector type,
     // there is direct shuffle support in LLVM. Use it!
     if (rank == 1 && v1Type == v2Type) {
-      Value shuffle = rewriter.create<LLVM::ShuffleVectorOp>(
+      Value llvmShuffleOp = rewriter.create<LLVM::ShuffleVectorOp>(
           loc, adaptor.v1(), adaptor.v2(), maskArrayAttr);
-      rewriter.replaceOp(op, shuffle);
+      rewriter.replaceOp(shuffleOp, llvmShuffleOp);
       return success();
     }
 
     // For all other cases, insert the individual values individually.
     Value insert = rewriter.create<LLVM::UndefOp>(loc, llvmType);
     int64_t insPos = 0;
     for (auto en : llvm::enumerate(maskArrayAttr)) {
       int64_t extPos = en.value().cast<IntegerAttr>().getInt();
       Value value = adaptor.v1();
       if (extPos >= v1Dim) {
         extPos -= v1Dim;
         value = adaptor.v2();
       }
       Value extract = extractOne(rewriter, *getTypeConverter(), loc, value,
                                  llvmType, rank, extPos);
       insert = insertOne(rewriter, *getTypeConverter(), loc, insert, extract,
                          llvmType, rank, insPos++);
     }
-    rewriter.replaceOp(op, insert);
+    rewriter.replaceOp(shuffleOp, insert);
     return success();
   }
 };
 
-class VectorExtractElementOpConversion : public ConvertToLLVMPattern {
+class VectorExtractElementOpConversion
+    : public ConvertOpToLLVMPattern<vector::ExtractElementOp> {
 public:
-  explicit VectorExtractElementOpConversion(MLIRContext *context,
-                                            LLVMTypeConverter &typeConverter)
-      : ConvertToLLVMPattern(vector::ExtractElementOp::getOperationName(),
-                             context, typeConverter) {}
+  using ConvertOpToLLVMPattern<
+      vector::ExtractElementOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(vector::ExtractElementOp extractEltOp,
+                  ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto adaptor = vector::ExtractElementOpAdaptor(operands);
-    auto extractEltOp = cast<vector::ExtractElementOp>(op);
     auto vectorType = extractEltOp.getVectorType();
     auto llvmType = typeConverter->convertType(vectorType.getElementType());
 
     // Bail if result type cannot be lowered.
     if (!llvmType)
       return failure();
 
     rewriter.replaceOpWithNewOp<LLVM::ExtractElementOp>(
-        op, llvmType, adaptor.vector(), adaptor.position());
+        extractEltOp, llvmType, adaptor.vector(), adaptor.position());
     return success();
   }
 };
 
-class VectorExtractOpConversion : public ConvertToLLVMPattern {
+class VectorExtractOpConversion
+    : public ConvertOpToLLVMPattern<vector::ExtractOp> {
 public:
-  explicit VectorExtractOpConversion(MLIRContext *context,
-                                     LLVMTypeConverter &typeConverter)
-      : ConvertToLLVMPattern(vector::ExtractOp::getOperationName(), context,
-                             typeConverter) {}
+  using ConvertOpToLLVMPattern<vector::ExtractOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(vector::ExtractOp extractOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
-    auto loc = op->getLoc();
+    auto loc = extractOp->getLoc();
     auto adaptor = vector::ExtractOpAdaptor(operands);
-    auto extractOp = cast<vector::ExtractOp>(op);
     auto vectorType = extractOp.getVectorType();
     auto resultType = extractOp.getResult().getType();
     auto llvmResultType = typeConverter->convertType(resultType);
     auto positionArrayAttr = extractOp.position();
 
     // Bail if result type cannot be lowered.
     if (!llvmResultType)
       return failure();
 
     // One-shot extraction of vector from array (only requires extractvalue).
     if (resultType.isa<VectorType>()) {
       Value extracted = rewriter.create<LLVM::ExtractValueOp>(
           loc, llvmResultType, adaptor.vector(), positionArrayAttr);
-      rewriter.replaceOp(op, extracted);
+      rewriter.replaceOp(extractOp, extracted);
       return success();
     }
 
     // Potential extraction of 1-D vector from array.
-    auto *context = op->getContext();
+    auto *context = extractOp->getContext();
     Value extracted = adaptor.vector();
     auto positionAttrs = positionArrayAttr.getValue();
     if (positionAttrs.size() > 1) {
       auto oneDVectorType = reducedVectorTypeBack(vectorType);
       auto nMinusOnePositionAttrs =
           ArrayAttr::get(positionAttrs.drop_back(), context);
       extracted = rewriter.create<LLVM::ExtractValueOp>(
           loc, typeConverter->convertType(oneDVectorType), extracted,
           nMinusOnePositionAttrs);
     }
 
     // Remaining extraction of element from 1-D LLVM vector
     auto position = positionAttrs.back().cast<IntegerAttr>();
     auto i64Type = LLVM::LLVMType::getInt64Ty(rewriter.getContext());
     auto constant = rewriter.create<LLVM::ConstantOp>(loc, i64Type, position);
     extracted =
         rewriter.create<LLVM::ExtractElementOp>(loc, extracted, constant);
-    rewriter.replaceOp(op, extracted);
+    rewriter.replaceOp(extractOp, extracted);
 
     return success();
   }
 };
 
 /// Conversion pattern that turns a vector.fma on a 1-D vector
 /// into an llvm.intr.fmuladd. This is a trivial 1-1 conversion.
 /// This does not match vectors of n >= 2 rank.
 ///
 /// Example:
 /// ```
 ///  vector.fma %a, %a, %a : vector<8xf32>
 /// ```
 /// is converted to:
 /// ```
 ///  llvm.intr.fmuladd %va, %va, %va:
 ///    (!llvm<"<8 x float>">, !llvm<"<8 x float>">, !llvm<"<8 x float>">)
 ///    -> !llvm<"<8 x float>">
 /// ```
-class VectorFMAOp1DConversion : public ConvertToLLVMPattern {
+class VectorFMAOp1DConversion : public ConvertOpToLLVMPattern<vector::FMAOp> {
 public:
-  explicit VectorFMAOp1DConversion(MLIRContext *context,
-                                   LLVMTypeConverter &typeConverter)
-      : ConvertToLLVMPattern(vector::FMAOp::getOperationName(), context,
-                             typeConverter) {}
+  using ConvertOpToLLVMPattern<vector::FMAOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(vector::FMAOp fmaOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto adaptor = vector::FMAOpAdaptor(operands);
-    vector::FMAOp fmaOp = cast<vector::FMAOp>(op);
     VectorType vType = fmaOp.getVectorType();
     if (vType.getRank() != 1)
       return failure();
-    rewriter.replaceOpWithNewOp<LLVM::FMulAddOp>(op, adaptor.lhs(),
+    rewriter.replaceOpWithNewOp<LLVM::FMulAddOp>(fmaOp, adaptor.lhs(),
                                                  adaptor.rhs(), adaptor.acc());
     return success();
   }
 };
 
-class VectorInsertElementOpConversion : public ConvertToLLVMPattern {
+class VectorInsertElementOpConversion
+    : public ConvertOpToLLVMPattern<vector::InsertElementOp> {
 public:
-  explicit VectorInsertElementOpConversion(MLIRContext *context,
-                                           LLVMTypeConverter &typeConverter)
-      : ConvertToLLVMPattern(vector::InsertElementOp::getOperationName(),
-                             context, typeConverter) {}
+  using ConvertOpToLLVMPattern<vector::InsertElementOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(vector::InsertElementOp insertEltOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto adaptor = vector::InsertElementOpAdaptor(operands);
-    auto insertEltOp = cast<vector::InsertElementOp>(op);
     auto vectorType = insertEltOp.getDestVectorType();
     auto llvmType = typeConverter->convertType(vectorType);
 
     // Bail if result type cannot be lowered.
     if (!llvmType)
       return failure();
 
     rewriter.replaceOpWithNewOp<LLVM::InsertElementOp>(
-        op, llvmType, adaptor.dest(), adaptor.source(), adaptor.position());
+        insertEltOp, llvmType, adaptor.dest(), adaptor.source(),
+        adaptor.position());
     return success();
   }
 };
 
-class VectorInsertOpConversion : public ConvertToLLVMPattern {
+class VectorInsertOpConversion
+    : public ConvertOpToLLVMPattern<vector::InsertOp> {
 public:
-  explicit VectorInsertOpConversion(MLIRContext *context,
-                                    LLVMTypeConverter &typeConverter)
-      : ConvertToLLVMPattern(vector::InsertOp::getOperationName(), context,
-                             typeConverter) {}
+  using ConvertOpToLLVMPattern<vector::InsertOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(vector::InsertOp insertOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
-    auto loc = op->getLoc();
+    auto loc = insertOp->getLoc();
     auto adaptor = vector::InsertOpAdaptor(operands);
-    auto insertOp = cast<vector::InsertOp>(op);
     auto sourceType = insertOp.getSourceType();
     auto destVectorType = insertOp.getDestVectorType();
     auto llvmResultType = typeConverter->convertType(destVectorType);
     auto positionArrayAttr = insertOp.position();
 
     // Bail if result type cannot be lowered.
     if (!llvmResultType)
       return failure();
 
     // One-shot insertion of a vector into an array (only requires insertvalue).
     if (sourceType.isa<VectorType>()) {
       Value inserted = rewriter.create<LLVM::InsertValueOp>(
           loc, llvmResultType, adaptor.dest(), adaptor.source(),
           positionArrayAttr);
-      rewriter.replaceOp(op, inserted);
+      rewriter.replaceOp(insertOp, inserted);
       return success();
     }
 
     // Potential extraction of 1-D vector from array.
-    auto *context = op->getContext();
+    auto *context = insertOp->getContext();
     Value extracted = adaptor.dest();
     auto positionAttrs = positionArrayAttr.getValue();
     auto position = positionAttrs.back().cast<IntegerAttr>();
     auto oneDVectorType = destVectorType;
     if (positionAttrs.size() > 1) {
       oneDVectorType = reducedVectorTypeBack(destVectorType);
       auto nMinusOnePositionAttrs =
           ArrayAttr::get(positionAttrs.drop_back(), context);
matchAndRewrite(vector::InsertOp insertOp, ArrayRef<Value> operands,
     if (positionAttrs.size() > 1) {
       auto nMinusOnePositionAttrs =
           ArrayAttr::get(positionAttrs.drop_back(), context);
       inserted = rewriter.create<LLVM::InsertValueOp>(loc, llvmResultType,
                                                       adaptor.dest(), inserted,
                                                       nMinusOnePositionAttrs);
     }
 
-    rewriter.replaceOp(op, inserted);
+    rewriter.replaceOp(insertOp, inserted);
     return success();
   }
 };
 
 /// Rank reducing rewrite for n-D FMA into (n-1)-D FMA where n > 1.
 ///
 /// Example:
 /// ```
if (ShapedType::isDynamic(sizes[index + 1]) ||
         ShapedType::isDynamicStrideOrOffset(strides[index + 1]))
       return None;
     if (strides[index] != strides[index + 1] * sizes[index + 1])
       return None;
   }
   return strides;
 }
 
-class VectorTypeCastOpConversion : public ConvertToLLVMPattern {
+class VectorTypeCastOpConversion
+    : public ConvertOpToLLVMPattern<vector::TypeCastOp> {
 public:
-  explicit VectorTypeCastOpConversion(MLIRContext *context,
-                                      LLVMTypeConverter &typeConverter)
-      : ConvertToLLVMPattern(vector::TypeCastOp::getOperationName(), context,
-                             typeConverter) {}
+  using ConvertOpToLLVMPattern<vector::TypeCastOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(vector::TypeCastOp castOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
-    auto loc = op->getLoc();
-    vector::TypeCastOp castOp = cast<vector::TypeCastOp>(op);
+    auto loc = castOp->getLoc();
     MemRefType sourceMemRefType =
         castOp.getOperand().getType().cast<MemRefType>();
     MemRefType targetMemRefType =
         castOp.getResult().getType().cast<MemRefType>();
 
     // Only static shape casts supported atm.
     if (!sourceMemRefType.hasStaticShape() ||
         !targetMemRefType.hasStaticShape())
for (auto indexedSize : llvm::enumerate(targetMemRefType.getShape())) {
       auto size = rewriter.create<LLVM::ConstantOp>(loc, int64Ty, sizeAttr);
       desc.setSize(rewriter, loc, index, size);
       auto strideAttr = rewriter.getIntegerAttr(rewriter.getIndexType(),
                                                 (*targetStrides)[index]);
       auto stride = rewriter.create<LLVM::ConstantOp>(loc, int64Ty, strideAttr);
       desc.setStride(rewriter, loc, index, stride);
     }
 
-    rewriter.replaceOp(op, {desc});
+    rewriter.replaceOp(castOp, {desc});
     return success();
   }
 };
 
 /// Conversion pattern that converts a 1-D vector transfer read/write op in a
 /// sequence of:
 /// 1. Get the source/dst address as an LLVM vector pointer.
 /// 2. Create a vector with linear indices [ 0 .. vector_length - 1 ].
 /// 3. Create an offsetVector = [ offset + 0 .. offset + vector_length - 1 ].
 /// 4. Create a mask where offsetVector is compared against memref upper bound.
 /// 5. Rewrite op as a masked read or write.
 template <typename ConcreteOp>
-class VectorTransferConversion : public ConvertToLLVMPattern {
+class VectorTransferConversion : public ConvertOpToLLVMPattern<ConcreteOp> {
 public:
-  explicit VectorTransferConversion(MLIRContext *context,
-                                    LLVMTypeConverter &typeConv,
+  explicit VectorTransferConversion(LLVMTypeConverter &typeConv,
                                     bool enableIndexOpt)
-      : ConvertToLLVMPattern(ConcreteOp::getOperationName(), context, typeConv),
+      : ConvertOpToLLVMPattern<ConcreteOp>(typeConv),
         enableIndexOptimizations(enableIndexOpt) {}
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(ConcreteOp xferOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
-    auto xferOp = cast<ConcreteOp>(op);
     auto adaptor = getTransferOpAdapter(xferOp, operands);
 
     if (xferOp.getVectorType().getRank() > 1 ||
         llvm::size(xferOp.indices()) == 0)
       return failure();
     if (xferOp.permutation_map() !=
         AffineMap::getMinorIdentityMap(xferOp.permutation_map().getNumInputs(),
                                        xferOp.getVectorType().getRank(),
-                                       op->getContext()))
+                                       xferOp->getContext()))
       return failure();
     // Only contiguous source tensors supported atm.
     auto strides = computeContiguousStrides(xferOp.getMemRefType());
     if (!strides)
       return failure();
 
-    auto toLLVMTy = [&](Type t) { return typeConverter->convertType(t); };
+    auto toLLVMTy = [&](Type t) {
+      return this->getTypeConverter()->convertType(t);
+    };
 
-    Location loc = op->getLoc();
+    Location loc = xferOp->getLoc();
     MemRefType memRefType = xferOp.getMemRefType();
 
     if (auto memrefVectorElementType =
             memRefType.getElementType().dyn_cast<VectorType>()) {
       // Memref has vector element type.
       if (memrefVectorElementType.getElementType() !=
           xferOp.getVectorType().getElementType())
         return failure();

 #endif // ifndef NDEBUG
     }
 
     // 1. Get the source/dst address as an LLVM vector pointer.
     //    The vector pointer would always be on address space 0, therefore
     //    addrspacecast shall be used when source/dst memrefs are not on
     //    address space 0.
     // TODO: support alignment when possible.
-    Value dataPtr = getStridedElementPtr(loc, memRefType, adaptor.memref(),
-                                         adaptor.indices(), rewriter);
+    Value dataPtr = this->getStridedElementPtr(
+        loc, memRefType, adaptor.memref(), adaptor.indices(), rewriter);
     auto vecTy =
         toLLVMTy(xferOp.getVectorType()).template cast<LLVM::LLVMType>();
     Value vectorDataPtr;
     if (memRefType.getMemorySpace() == 0)
       vectorDataPtr =
           rewriter.create<LLVM::BitcastOp>(loc, vecTy.getPointerTo(), dataPtr);
     else
       vectorDataPtr = rewriter.create<LLVM::AddrSpaceCastOp>(
           loc, vecTy.getPointerTo(), dataPtr);
 
     if (!xferOp.isMaskedDim(0))
-      return replaceTransferOpWithLoadOrStore(
-          rewriter, *getTypeConverter(), loc, xferOp, operands, vectorDataPtr);
+      return replaceTransferOpWithLoadOrStore(rewriter,
+                                              *this->getTypeConverter(), loc,
+                                              xferOp, operands, vectorDataPtr);
 
     // 2. Create a vector with linear indices [ 0 .. vector_length - 1 ].
     // 3. Create offsetVector = [ offset + 0 .. offset + vector_length - 1 ].
     // 4. Let dim the memref dimension, compute the vector comparison mask:
     //   [ offset + 0 .. offset + vector_length - 1 ] < [ dim .. dim ]
     //
     // TODO: when the leaf transfer rank is k > 1, we need the last `k`
     //       dimensions here.
     unsigned vecWidth = vecTy.getVectorNumElements();
     unsigned lastIndex = llvm::size(xferOp.indices()) - 1;
     Value off = xferOp.indices()[lastIndex];
     Value dim = rewriter.create<DimOp>(loc, xferOp.memref(), lastIndex);
-    Value mask = buildVectorComparison(rewriter, op, enableIndexOptimizations,
-                                       vecWidth, dim, &off);
+    Value mask = buildVectorComparison(
+        rewriter, xferOp, enableIndexOptimizations, vecWidth, dim, &off);
 
     // 5. Rewrite as a masked read / write.
-    return replaceTransferOpWithMasked(rewriter, *getTypeConverter(), loc,
+    return replaceTransferOpWithMasked(rewriter, *this->getTypeConverter(), loc,
                                        xferOp, operands, vectorDataPtr, mask);
   }
 
 private:
   const bool enableIndexOptimizations;
 };
 
-class VectorPrintOpConversion : public ConvertToLLVMPattern {
+class VectorPrintOpConversion : public ConvertOpToLLVMPattern<vector::PrintOp> {
 public:
-  explicit VectorPrintOpConversion(MLIRContext *context,
-                                   LLVMTypeConverter &typeConverter)
-      : ConvertToLLVMPattern(vector::PrintOp::getOperationName(), context,
-                             typeConverter) {}
+  using ConvertOpToLLVMPattern<vector::PrintOp>::ConvertOpToLLVMPattern;
 
   // Proof-of-concept lowering implementation that relies on a small
   // runtime support library, which only needs to provide a few
   // printing methods (single value for all data types, opening/closing
   // bracket, comma, newline). The lowering fully unrolls a vector
   // in terms of these elementary printing operations. The advantage
   // of this approach is that the library can remain unaware of all
   // low-level implementation details of vectors while still supporting
   // output of any shaped and dimensioned vector. Due to full unrolling,
   // this approach is less suited for very large vectors though.
   //
   // TODO: rely solely on libc in future? something else?
   //
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(vector::PrintOp printOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
-    auto printOp = cast<vector::PrintOp>(op);
     auto adaptor = vector::PrintOpAdaptor(operands);
     Type printType = printOp.getPrintType();
 
     if (typeConverter->convertType(printType) == nullptr)
       return failure();
 
     // Make sure element type has runtime support.
     PrintConversion conversion = PrintConversion::None;
     VectorType vectorType = printType.dyn_cast<VectorType>();
     Type eltType = vectorType ? vectorType.getElementType() : printType;
     Operation *printer;
     if (eltType.isF32()) {
-      printer = getPrintFloat(op);
+      printer = getPrintFloat(printOp);
     } else if (eltType.isF64()) {
-      printer = getPrintDouble(op);
+      printer = getPrintDouble(printOp);
     } else if (eltType.isIndex()) {
-      printer = getPrintU64(op);
+      printer = getPrintU64(printOp);
     } else if (auto intTy = eltType.dyn_cast<IntegerType>()) {
       // Integers need a zero or sign extension on the operand
       // (depending on the source type) as well as a signed or
       // unsigned print method. Up to 64-bit is supported.
       unsigned width = intTy.getWidth();
       if (intTy.isUnsigned()) {
         if (width <= 64) {
           if (width < 64)
             conversion = PrintConversion::ZeroExt64;
-          printer = getPrintU64(op);
+          printer = getPrintU64(printOp);
         } else {
           return failure();
         }
       } else {
         assert(intTy.isSignless() || intTy.isSigned());
         if (width <= 64) {
           // Note that we *always* zero extend booleans (1-bit integers),
           // so that true/false is printed as 1/0 rather than -1/0.
           if (width == 1)
             conversion = PrintConversion::ZeroExt64;
           else if (width < 64)
             conversion = PrintConversion::SignExt64;
-          printer = getPrintI64(op);
+          printer = getPrintI64(printOp);
         } else {
           return failure();
         }
       }
     } else {
       return failure();
     }
 
     // Unroll vector into elementary print calls.
     int64_t rank = vectorType ? vectorType.getRank() : 0;
-    emitRanks(rewriter, op, adaptor.source(), vectorType, printer, rank,
+    emitRanks(rewriter, printOp, adaptor.source(), vectorType, printer, rank,
               conversion);
-    emitCall(rewriter, op->getLoc(), getPrintNewline(op));
-    rewriter.eraseOp(op);
+    emitCall(rewriter, printOp->getLoc(), getPrintNewline(printOp));
+    rewriter.eraseOp(printOp);
     return success();
   }
 
 private:
   enum class PrintConversion {
     // clang-format off
     None,
     ZeroExt64,
void mlir::populateVectorToLLVMConversionPatterns(
     bool reassociateFPReductions, bool enableIndexOptimizations) {
   MLIRContext *ctx = converter.getDialect()->getContext();
   // clang-format off
   patterns.insert<VectorFMAOpNDRewritePattern,
                   VectorInsertStridedSliceOpDifferentRankRewritePattern,
                   VectorInsertStridedSliceOpSameRankRewritePattern,
                   VectorExtractStridedSliceOpConversion>(ctx);
   patterns.insert<VectorReductionOpConversion>(
-      ctx, converter, reassociateFPReductions);
+      converter, reassociateFPReductions);
   patterns.insert<VectorCreateMaskOpConversion,
                   VectorTransferConversion<TransferReadOp>,
                   VectorTransferConversion<TransferWriteOp>>(
-      ctx, converter, enableIndexOptimizations);
+      converter, enableIndexOptimizations);
   patterns
       .insert<VectorShuffleOpConversion,
               VectorExtractElementOpConversion,
               VectorExtractOpConversion,
               VectorFMAOp1DConversion,
               VectorInsertElementOpConversion,
               VectorInsertOpConversion,
               VectorPrintOpConversion,
               VectorTypeCastOpConversion,
               VectorMaskedLoadOpConversion,
               VectorMaskedStoreOpConversion,
               VectorGatherOpConversion,
               VectorScatterOpConversion,
               VectorExpandLoadOpConversion,
-              VectorCompressStoreOpConversion>(ctx, converter);
+              VectorCompressStoreOpConversion>(converter);
   // clang-format on
 }
 
 void mlir::populateVectorToLLVMMatrixConversionPatterns(
     LLVMTypeConverter &converter, OwningRewritePatternList &patterns) {
-  MLIRContext *ctx = converter.getDialect()->getContext();
-  patterns.insert<VectorMatmulOpConversion>(ctx, converter);
-  patterns.insert<VectorFlatTransposeOpConversion>(ctx, converter);
+  patterns.insert<VectorMatmulOpConversion>(converter);
+  patterns.insert<VectorFlatTransposeOpConversion>(converter);
 }

mlir/lib/Conversion/VectorToROCDL/VectorToROCDL.cpp

 }
 
 namespace {
 /// Conversion pattern that converts a 1-D vector transfer read/write.
 /// Note that this conversion pass only converts vector x2 or x4 f32
 /// types. For unsupported cases, they will fall back to the vector to
 /// llvm conversion pattern.
 template <typename ConcreteOp>
-class VectorTransferConversion : public ConvertToLLVMPattern {
+class VectorTransferConversion : public ConvertOpToLLVMPattern<ConcreteOp> {
 public:
-  explicit VectorTransferConversion(MLIRContext *context,
-                                    LLVMTypeConverter &typeConv)
-      : ConvertToLLVMPattern(ConcreteOp::getOperationName(), context,
-                             typeConv) {}
+  using ConvertOpToLLVMPattern<ConcreteOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+  matchAndRewrite(ConcreteOp xferOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
-    auto xferOp = cast<ConcreteOp>(op);
     typename ConcreteOp::Adaptor adaptor(operands);
 
     if (xferOp.getVectorType().getRank() > 1 ||
         llvm::size(xferOp.indices()) == 0)
       return failure();
 
     if (!xferOp.permutation_map().isMinorIdentity())
       return failure();
 
     // Have it handled in vector->llvm conversion pass.
     if (!xferOp.isMaskedDim(0))
       return failure();
 
-    auto toLLVMTy = [&](Type t) { return typeConverter->convertType(t); };
+    auto toLLVMTy = [&](Type t) {
+      return this->getTypeConverter()->convertType(t);
+    };
     LLVM::LLVMType vecTy =
         toLLVMTy(xferOp.getVectorType()).template cast<LLVM::LLVMType>();
     unsigned vecWidth = vecTy.getVectorNumElements();
-    Location loc = op->getLoc();
+    Location loc = xferOp->getLoc();
 
     // The backend result vector scalarization have trouble scalarize
     // <1 x ty> result, exclude the x1 width from the lowering.
     if (vecWidth != 2 && vecWidth != 4)
       return failure();
 
     // Obtain dataPtr and elementType from the memref.
     MemRefType memRefType = xferOp.getMemRefType();
     // MUBUF instruction operate only on addresspace 0(unified) or 1(global)
     // In case of 3(LDS): fall back to vector->llvm pass
     // In case of 5(VGPR): wrong
     if ((memRefType.getMemorySpace() != 0) &&
         (memRefType.getMemorySpace() != 1))
       return failure();
 
     // Note that the dataPtr starts at the offset address specified by
     // indices, so no need to calculate offset size in bytes again in
     // the MUBUF instruction.
-    Value dataPtr = getStridedElementPtr(loc, memRefType, adaptor.memref(),
-                                         adaptor.indices(), rewriter);
+    Value dataPtr = this->getStridedElementPtr(
+        loc, memRefType, adaptor.memref(), adaptor.indices(), rewriter);
 
     // 1. Create and fill a <4 x i32> dwordConfig with:
     //    1st two elements holding the address of dataPtr.
     //    3rd element: -1.
     //    4th element: 0x27000.
     SmallVector<int32_t, 4> constConfigAttr{0, 0, -1, 0x27000};
     Type i32Ty = rewriter.getIntegerType(32);
     VectorType i32Vecx4 = VectorType::get(4, i32Ty);
     Value constConfig = rewriter.create<LLVM::ConstantOp>(
         loc, toLLVMTy(i32Vecx4),
         DenseElementsAttr::get(i32Vecx4, ArrayRef<int32_t>(constConfigAttr)));
 
     // Treat first two element of <4 x i32> as i64, and save the dataPtr
     // to it.
     Type i64Ty = rewriter.getIntegerType(64);
     Value i64x2Ty = rewriter.create<LLVM::BitcastOp>(
         loc,
         LLVM::LLVMType::getVectorTy(
             toLLVMTy(i64Ty).template cast<LLVM::LLVMType>(), 2),
         constConfig);
     Value dataPtrAsI64 = rewriter.create<LLVM::PtrToIntOp>(
         loc, toLLVMTy(i64Ty).template cast<LLVM::LLVMType>(), dataPtr);
-    Value zero = createIndexConstant(rewriter, loc, 0);
+    Value zero = this->createIndexConstant(rewriter, loc, 0);
     Value dwordConfig = rewriter.create<LLVM::InsertElementOp>(
         loc,
         LLVM::LLVMType::getVectorTy(
             toLLVMTy(i64Ty).template cast<LLVM::LLVMType>(), 2),
         i64x2Ty, dataPtrAsI64, zero);
     dwordConfig =
         rewriter.create<LLVM::BitcastOp>(loc, toLLVMTy(i32Vecx4), dwordConfig);
 
     // 2. Rewrite op as a buffer read or write.
     Value int1False = rewriter.create<LLVM::ConstantOp>(
         loc, toLLVMTy(rewriter.getIntegerType(1)),
         rewriter.getIntegerAttr(rewriter.getIntegerType(1), 0));
     Value int32Zero = rewriter.create<LLVM::ConstantOp>(
         loc, toLLVMTy(i32Ty),
         rewriter.getIntegerAttr(rewriter.getIntegerType(32), 0));
     return replaceTransferOpWithMubuf(
-        rewriter, operands, *getTypeConverter(), loc, xferOp, vecTy,
+        rewriter, operands, *this->getTypeConverter(), loc, xferOp, vecTy,
         dwordConfig, int32Zero, int32Zero, int1False, int1False);
   }
 };
 } // end anonymous namespace
 
 void mlir::populateVectorToROCDLConversionPatterns(
     LLVMTypeConverter &converter, OwningRewritePatternList &patterns) {
-  MLIRContext *ctx = converter.getDialect()->getContext();
   patterns.insert<VectorTransferConversion<TransferReadOp>,
-                  VectorTransferConversion<TransferWriteOp>>(ctx, converter);
+                  VectorTransferConversion<TransferWriteOp>>(converter);
 }
 
 namespace {
 struct LowerVectorToROCDLPass
     : public ConvertVectorToROCDLBase<LowerVectorToROCDLPass> {
   void runOnOperation() override;
 };
 } // namespace