[mlir:OpConversion] Remove the remaing usages of the deprecated matchAndRewrite methods

[rriddle]

Summary

This commits updates the remaining usages of the ArrayRef<Value> based
matchAndRewrite/rewrite methods in favor of the new OpAdaptor
overload.

Depends On D110293

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

 class OneToOneConvertToLLVMPattern : public ConvertOpToLLVMPattern<SourceOp> {
 public:
   using ConvertOpToLLVMPattern<SourceOp>::ConvertOpToLLVMPattern;
   using Super = OneToOneConvertToLLVMPattern<SourceOp, TargetOp>;
 
   /// Converts the type of the result to an LLVM type, pass operands as is,
   /// preserve attributes.
   LogicalResult
-  matchAndRewrite(SourceOp op, ArrayRef<Value> operands,
+  matchAndRewrite(SourceOp op, typename SourceOp::Adaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     return LLVM::detail::oneToOneRewrite(op, TargetOp::getOperationName(),
-                                         operands, *this->getTypeConverter(),
-                                         rewriter);
+                                         adaptor.getOperands(),
+                                         *this->getTypeConverter(), rewriter);
   }
 };
 
 } // namespace mlir
 
 #endif // MLIR_CONVERSION_LLVMCOMMON_PATTERN_H

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

 /// LLVM Dialect. This supports higher-dimensional vector types.
 template <typename SourceOp, typename TargetOp>
 class VectorConvertToLLVMPattern : public ConvertOpToLLVMPattern<SourceOp> {
 public:
   using ConvertOpToLLVMPattern<SourceOp>::ConvertOpToLLVMPattern;
   using Super = VectorConvertToLLVMPattern<SourceOp, TargetOp>;
 
   LogicalResult
-  matchAndRewrite(SourceOp op, ArrayRef<Value> operands,
+  matchAndRewrite(SourceOp op, typename SourceOp::Adaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     static_assert(
         std::is_base_of<OpTrait::OneResult<SourceOp>, SourceOp>::value,
         "expected single result op");
     return LLVM::detail::vectorOneToOneRewrite(
-        op, TargetOp::getOperationName(), operands, *this->getTypeConverter(),
-        rewriter);
+        op, TargetOp::getOperationName(), adaptor.getOperands(),
+        *this->getTypeConverter(), rewriter);
   }
 };
 } // namespace mlir
 
 #endif // MLIR_CONVERSION_LLVMCOMMON_VECTORPATTERN_H

mlir/lib/Conversion/ComplexToLLVM/ComplexToLLVM.cpp

 //===----------------------------------------------------------------------===//
 
 namespace {
 
 struct AbsOpConversion : public ConvertOpToLLVMPattern<complex::AbsOp> {
   using ConvertOpToLLVMPattern<complex::AbsOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(complex::AbsOp op, ArrayRef<Value> operands,
+  matchAndRewrite(complex::AbsOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    complex::AbsOp::Adaptor transformed(operands);
     auto loc = op.getLoc();
 
-    ComplexStructBuilder complexStruct(transformed.complex());
+    ComplexStructBuilder complexStruct(adaptor.complex());
     Value real = complexStruct.real(rewriter, op.getLoc());
     Value imag = complexStruct.imaginary(rewriter, op.getLoc());
 
     auto fmf = LLVM::FMFAttr::get(op.getContext(), {});
     Value sqNorm = rewriter.create<LLVM::FAddOp>(
         loc, rewriter.create<LLVM::FMulOp>(loc, real, real, fmf),
         rewriter.create<LLVM::FMulOp>(loc, imag, imag, fmf), fmf);
 
     rewriter.replaceOpWithNewOp<LLVM::SqrtOp>(op, sqNorm);
     return success();
   }
 };
 
 struct CreateOpConversion : public ConvertOpToLLVMPattern<complex::CreateOp> {
   using ConvertOpToLLVMPattern<complex::CreateOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(complex::CreateOp complexOp, ArrayRef<Value> operands,
+  matchAndRewrite(complex::CreateOp complexOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    complex::CreateOp::Adaptor transformed(operands);
-
     // Pack real and imaginary part in a complex number struct.
     auto loc = complexOp.getLoc();
     auto structType = typeConverter->convertType(complexOp.getType());
     auto complexStruct = ComplexStructBuilder::undef(rewriter, loc, structType);
-    complexStruct.setReal(rewriter, loc, transformed.real());
-    complexStruct.setImaginary(rewriter, loc, transformed.imaginary());
+    complexStruct.setReal(rewriter, loc, adaptor.real());
+    complexStruct.setImaginary(rewriter, loc, adaptor.imaginary());
 
     rewriter.replaceOp(complexOp, {complexStruct});
     return success();
   }
 };
 
 struct ReOpConversion : public ConvertOpToLLVMPattern<complex::ReOp> {
   using ConvertOpToLLVMPattern<complex::ReOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(complex::ReOp op, ArrayRef<Value> operands,
+  matchAndRewrite(complex::ReOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    complex::ReOp::Adaptor transformed(operands);
-
     // Extract real part from the complex number struct.
-    ComplexStructBuilder complexStruct(transformed.complex());
+    ComplexStructBuilder complexStruct(adaptor.complex());
     Value real = complexStruct.real(rewriter, op.getLoc());
     rewriter.replaceOp(op, real);
 
     return success();
   }
 };
 
 struct ImOpConversion : public ConvertOpToLLVMPattern<complex::ImOp> {
   using ConvertOpToLLVMPattern<complex::ImOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(complex::ImOp op, ArrayRef<Value> operands,
+  matchAndRewrite(complex::ImOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    complex::ImOp::Adaptor transformed(operands);
-
     // Extract imaginary part from the complex number struct.
-    ComplexStructBuilder complexStruct(transformed.complex());
+    ComplexStructBuilder complexStruct(adaptor.complex());
     Value imaginary = complexStruct.imaginary(rewriter, op.getLoc());
     rewriter.replaceOp(op, imaginary);
 
     return success();
   }
 };
 
 struct BinaryComplexOperands {
   std::complex<Value> lhs;
   std::complex<Value> rhs;
 };
 
 template <typename OpTy>
 BinaryComplexOperands
-unpackBinaryComplexOperands(OpTy op, ArrayRef<Value> operands,
+unpackBinaryComplexOperands(OpTy op, typename OpTy::Adaptor adaptor,
                             ConversionPatternRewriter &rewriter) {
   auto loc = op.getLoc();
-  typename OpTy::Adaptor transformed(operands);
 
   // Extract real and imaginary values from operands.
   BinaryComplexOperands unpacked;
-  ComplexStructBuilder lhs(transformed.lhs());
+  ComplexStructBuilder lhs(adaptor.lhs());
   unpacked.lhs.real(lhs.real(rewriter, loc));
   unpacked.lhs.imag(lhs.imaginary(rewriter, loc));
-  ComplexStructBuilder rhs(transformed.rhs());
+  ComplexStructBuilder rhs(adaptor.rhs());
   unpacked.rhs.real(rhs.real(rewriter, loc));
   unpacked.rhs.imag(rhs.imaginary(rewriter, loc));
 
   return unpacked;
 }
 
 struct AddOpConversion : public ConvertOpToLLVMPattern<complex::AddOp> {
   using ConvertOpToLLVMPattern<complex::AddOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(complex::AddOp op, ArrayRef<Value> operands,
+  matchAndRewrite(complex::AddOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     auto loc = op.getLoc();
     BinaryComplexOperands arg =
-        unpackBinaryComplexOperands<complex::AddOp>(op, operands, rewriter);
+        unpackBinaryComplexOperands<complex::AddOp>(op, adaptor, rewriter);
 
     // Initialize complex number struct for result.
     auto structType = typeConverter->convertType(op.getType());
     auto result = ComplexStructBuilder::undef(rewriter, loc, structType);
 
     // Emit IR to add complex numbers.
     auto fmf = LLVM::FMFAttr::get(op.getContext(), {});
     Value real =
         rewriter.create<LLVM::FAddOp>(loc, arg.lhs.real(), arg.rhs.real(), fmf);
     Value imag =
         rewriter.create<LLVM::FAddOp>(loc, arg.lhs.imag(), arg.rhs.imag(), fmf);
     result.setReal(rewriter, loc, real);
     result.setImaginary(rewriter, loc, imag);
 
     rewriter.replaceOp(op, {result});
     return success();
   }
 };
 
 struct DivOpConversion : public ConvertOpToLLVMPattern<complex::DivOp> {
   using ConvertOpToLLVMPattern<complex::DivOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(complex::DivOp op, ArrayRef<Value> operands,
+  matchAndRewrite(complex::DivOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     auto loc = op.getLoc();
     BinaryComplexOperands arg =
-        unpackBinaryComplexOperands<complex::DivOp>(op, operands, rewriter);
+        unpackBinaryComplexOperands<complex::DivOp>(op, adaptor, rewriter);
 
     // Initialize complex number struct for result.
     auto structType = typeConverter->convertType(op.getType());
     auto result = ComplexStructBuilder::undef(rewriter, loc, structType);
 
     // Emit IR to add complex numbers.
     auto fmf = LLVM::FMFAttr::get(op.getContext(), {});
     Value rhsRe = arg.rhs.real();
matchAndRewrite(complex::DivOp op, OpAdaptor adaptor,
     return success();
   }
 };
 
 struct MulOpConversion : public ConvertOpToLLVMPattern<complex::MulOp> {
   using ConvertOpToLLVMPattern<complex::MulOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(complex::MulOp op, ArrayRef<Value> operands,
+  matchAndRewrite(complex::MulOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     auto loc = op.getLoc();
     BinaryComplexOperands arg =
-        unpackBinaryComplexOperands<complex::MulOp>(op, operands, rewriter);
+        unpackBinaryComplexOperands<complex::MulOp>(op, adaptor, rewriter);
 
     // Initialize complex number struct for result.
     auto structType = typeConverter->convertType(op.getType());
     auto result = ComplexStructBuilder::undef(rewriter, loc, structType);
 
     // Emit IR to add complex numbers.
     auto fmf = LLVM::FMFAttr::get(op.getContext(), {});
     Value rhsRe = arg.rhs.real();
matchAndRewrite(complex::MulOp op, OpAdaptor adaptor,
     return success();
   }
 };
 
 struct SubOpConversion : public ConvertOpToLLVMPattern<complex::SubOp> {
   using ConvertOpToLLVMPattern<complex::SubOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(complex::SubOp op, ArrayRef<Value> operands,
+  matchAndRewrite(complex::SubOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     auto loc = op.getLoc();
     BinaryComplexOperands arg =
-        unpackBinaryComplexOperands<complex::SubOp>(op, operands, rewriter);
+        unpackBinaryComplexOperands<complex::SubOp>(op, adaptor, rewriter);
 
     // Initialize complex number struct for result.
     auto structType = typeConverter->convertType(op.getType());
     auto result = ComplexStructBuilder::undef(rewriter, loc, structType);
 
     // Emit IR to substract complex numbers.
     auto fmf = LLVM::FMFAttr::get(op.getContext(), {});
     Value real =

mlir/lib/Conversion/GPUCommon/GPUOpsLowering.h

 struct GPUFuncOpLowering : ConvertOpToLLVMPattern<gpu::GPUFuncOp> {
   GPUFuncOpLowering(LLVMTypeConverter &converter, unsigned allocaAddrSpace,
                     Identifier kernelAttributeName)
       : ConvertOpToLLVMPattern<gpu::GPUFuncOp>(converter),
         allocaAddrSpace(allocaAddrSpace),
         kernelAttributeName(kernelAttributeName) {}
 
   LogicalResult
-  matchAndRewrite(gpu::GPUFuncOp gpuFuncOp, ArrayRef<Value> operands,
+  matchAndRewrite(gpu::GPUFuncOp gpuFuncOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override;
 
 private:
   /// The address spcae to use for `alloca`s in private memory.
   unsigned allocaAddrSpace;
 
   /// The attribute name to use instead of `gpu.kernel`.
   Identifier kernelAttributeName;
 };
 
 struct GPUReturnOpLowering : public ConvertOpToLLVMPattern<gpu::ReturnOp> {
   using ConvertOpToLLVMPattern<gpu::ReturnOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(gpu::ReturnOp op, ArrayRef<Value> operands,
+  matchAndRewrite(gpu::ReturnOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    rewriter.replaceOpWithNewOp<LLVM::ReturnOp>(op, operands);
+    rewriter.replaceOpWithNewOp<LLVM::ReturnOp>(op, adaptor.getOperands());
     return success();
   }
 };
 
 } // namespace mlir
 
 #endif // MLIR_CONVERSION_GPUCOMMON_GPUOPSLOWERING_H_

mlir/lib/Conversion/GPUCommon/GPUOpsLowering.cpp

 //===- GPUOpsLowering.cpp - GPU FuncOp / ReturnOp lowering ----------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #include "GPUOpsLowering.h"
 #include "mlir/Dialect/StandardOps/IR/Ops.h"
 #include "mlir/IR/Builders.h"
 #include "llvm/Support/FormatVariadic.h"
 
 using namespace mlir;
 
 LogicalResult
-GPUFuncOpLowering::matchAndRewrite(gpu::GPUFuncOp gpuFuncOp,
-                                   ArrayRef<Value> operands,
+GPUFuncOpLowering::matchAndRewrite(gpu::GPUFuncOp gpuFuncOp, OpAdaptor adaptor,
                                    ConversionPatternRewriter &rewriter) const {
-  assert(operands.empty() && "func op is not expected to have operands");
   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>();

mlir/lib/Conversion/GPUCommon/GPUToLLVMConversion.cpp

 class ConvertHostRegisterOpToGpuRuntimeCallPattern
     : public ConvertOpToGpuRuntimeCallPattern<gpu::HostRegisterOp> {
 public:
   ConvertHostRegisterOpToGpuRuntimeCallPattern(LLVMTypeConverter &typeConverter)
       : ConvertOpToGpuRuntimeCallPattern<gpu::HostRegisterOp>(typeConverter) {}
 
 private:
   LogicalResult
-  matchAndRewrite(gpu::HostRegisterOp hostRegisterOp, ArrayRef<Value> operands,
+  matchAndRewrite(gpu::HostRegisterOp hostRegisterOp, 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)
       : ConvertOpToGpuRuntimeCallPattern<gpu::AllocOp>(typeConverter) {}
 
 private:
   LogicalResult
-  matchAndRewrite(gpu::AllocOp allocOp, ArrayRef<Value> operands,
+  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)
       : ConvertOpToGpuRuntimeCallPattern<gpu::DeallocOp>(typeConverter) {}
 
 private:
   LogicalResult
-  matchAndRewrite(gpu::DeallocOp deallocOp, ArrayRef<Value> operands,
+  matchAndRewrite(gpu::DeallocOp deallocOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
 class ConvertAsyncYieldToGpuRuntimeCallPattern
     : public ConvertOpToGpuRuntimeCallPattern<async::YieldOp> {
 public:
   ConvertAsyncYieldToGpuRuntimeCallPattern(LLVMTypeConverter &typeConverter)
       : ConvertOpToGpuRuntimeCallPattern<async::YieldOp>(typeConverter) {}
 
 private:
   LogicalResult
-  matchAndRewrite(async::YieldOp yieldOp, ArrayRef<Value> operands,
+  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)
       : ConvertOpToGpuRuntimeCallPattern<gpu::WaitOp>(typeConverter) {}
 
 private:
   LogicalResult
-  matchAndRewrite(gpu::WaitOp waitOp, ArrayRef<Value> operands,
+  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)
       : ConvertOpToGpuRuntimeCallPattern<gpu::WaitOp>(typeConverter) {}
 
 private:
   LogicalResult
-  matchAndRewrite(gpu::WaitOp waitOp, ArrayRef<Value> operands,
+  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:
class ConvertLaunchFuncOpToGpuRuntimeCallPattern
     : public ConvertOpToGpuRuntimeCallPattern<gpu::LaunchFuncOp> {
 public:
   ConvertLaunchFuncOpToGpuRuntimeCallPattern(LLVMTypeConverter &typeConverter,
                                              StringRef gpuBinaryAnnotation)
       : ConvertOpToGpuRuntimeCallPattern<gpu::LaunchFuncOp>(typeConverter),
         gpuBinaryAnnotation(gpuBinaryAnnotation) {}
 
 private:
-  Value generateParamsArray(gpu::LaunchFuncOp launchOp,
-                            ArrayRef<Value> operands, OpBuilder &builder) const;
+  Value generateParamsArray(gpu::LaunchFuncOp launchOp, OpAdaptor adaptor,
+                            OpBuilder &builder) const;
   Value generateKernelNameConstant(StringRef moduleName, StringRef name,
                                    Location loc, OpBuilder &builder) const;
 
   LogicalResult
-  matchAndRewrite(gpu::LaunchFuncOp launchOp, ArrayRef<Value> operands,
+  matchAndRewrite(gpu::LaunchFuncOp launchOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override;
 
   llvm::SmallString<32> gpuBinaryAnnotation;
 };
 
 class EraseGpuModuleOpPattern : public OpRewritePattern<gpu::GPUModuleOp> {
   using OpRewritePattern<gpu::GPUModuleOp>::OpRewritePattern;
 

 class ConvertMemcpyOpToGpuRuntimeCallPattern
     : public ConvertOpToGpuRuntimeCallPattern<gpu::MemcpyOp> {
 public:
   ConvertMemcpyOpToGpuRuntimeCallPattern(LLVMTypeConverter &typeConverter)
       : ConvertOpToGpuRuntimeCallPattern<gpu::MemcpyOp>(typeConverter) {}
 
 private:
   LogicalResult
-  matchAndRewrite(gpu::MemcpyOp memcpyOp, ArrayRef<Value> operands,
+  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)
       : ConvertOpToGpuRuntimeCallPattern<gpu::MemsetOp>(typeConverter) {}
 
 private:
   LogicalResult
-  matchAndRewrite(gpu::MemsetOp memsetOp, ArrayRef<Value> operands,
+  matchAndRewrite(gpu::MemsetOp memsetOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override;
 };
 } // namespace
 
 void GpuToLLVMConversionPass::runOnOperation() {
   LLVMTypeConverter converter(&getContext());
   RewritePatternSet patterns(&getContext());
   LLVMConversionTarget target(getContext());
isAsyncWithOneDependency(ConversionPatternRewriter &rewriter,
 
   if (!op.getAsyncToken())
     return rewriter.notifyMatchFailure(op, "Can convert only async version.");
 
   return success();
 }
 
 LogicalResult ConvertHostRegisterOpToGpuRuntimeCallPattern::matchAndRewrite(
-    gpu::HostRegisterOp hostRegisterOp, ArrayRef<Value> operands,
+    gpu::HostRegisterOp hostRegisterOp, OpAdaptor adaptor,
     ConversionPatternRewriter &rewriter) const {
   auto *op = hostRegisterOp.getOperation();
-  if (failed(areAllLLVMTypes(op, operands, rewriter)))
+  if (failed(areAllLLVMTypes(op, adaptor.getOperands(), rewriter)))
     return failure();
 
   Location loc = op->getLoc();
 
   auto memRefType = hostRegisterOp.value().getType();
   auto elementType = memRefType.cast<UnrankedMemRefType>().getElementType();
   auto elementSize = getSizeInBytes(loc, elementType, rewriter);
 
-  auto arguments = getTypeConverter()->promoteOperands(loc, op->getOperands(),
-                                                       operands, rewriter);
+  auto arguments = getTypeConverter()->promoteOperands(
+      loc, op->getOperands(), adaptor.getOperands(), rewriter);
   arguments.push_back(elementSize);
   hostRegisterCallBuilder.create(loc, rewriter, arguments);
 
   rewriter.eraseOp(op);
   return success();
 }
 
 LogicalResult ConvertAllocOpToGpuRuntimeCallPattern::matchAndRewrite(
-    gpu::AllocOp allocOp, ArrayRef<Value> operands,
+    gpu::AllocOp allocOp, OpAdaptor adaptor,
     ConversionPatternRewriter &rewriter) const {
   MemRefType memRefType = allocOp.getType();
 
-  if (failed(areAllLLVMTypes(allocOp, operands, rewriter)) ||
+  if (failed(areAllLLVMTypes(allocOp, adaptor.getOperands(), rewriter)) ||
       !isConvertibleAndHasIdentityMaps(memRefType) ||
       failed(isAsyncWithOneDependency(rewriter, allocOp)))
     return failure();
 
   auto loc = allocOp.getLoc();
-  auto adaptor = gpu::AllocOpAdaptor(operands, allocOp->getAttrDictionary());
 
   // Get shape of the memref as values: static sizes are constant
   // values and dynamic sizes are passed to 'alloc' as operands.
   SmallVector<Value, 4> shape;
   SmallVector<Value, 4> strides;
   Value sizeBytes;
   getMemRefDescriptorSizes(loc, memRefType, adaptor.dynamicSizes(), rewriter,
                            shape, strides, sizeBytes);
auto memRefDescriptor = this->createMemRefDescriptor(
       loc, memRefType, allocatedPtr, alignedPtr, shape, strides, rewriter);
 
   rewriter.replaceOp(allocOp, {memRefDescriptor, stream});
 
   return success();
 }
 
 LogicalResult ConvertDeallocOpToGpuRuntimeCallPattern::matchAndRewrite(
-    gpu::DeallocOp deallocOp, ArrayRef<Value> operands,
+    gpu::DeallocOp deallocOp, OpAdaptor adaptor,
     ConversionPatternRewriter &rewriter) const {
-  if (failed(areAllLLVMTypes(deallocOp, operands, rewriter)) ||
+  if (failed(areAllLLVMTypes(deallocOp, adaptor.getOperands(), rewriter)) ||
       failed(isAsyncWithOneDependency(rewriter, deallocOp)))
     return failure();
 
   Location loc = deallocOp.getLoc();
 
-  auto adaptor =
-      gpu::DeallocOpAdaptor(operands, deallocOp->getAttrDictionary());
   Value pointer =
       MemRefDescriptor(adaptor.memref()).allocatedPtr(rewriter, loc);
   auto casted = rewriter.create<LLVM::BitcastOp>(loc, llvmPointerType, pointer);
   Value stream = adaptor.asyncDependencies().front();
   deallocCallBuilder.create(loc, rewriter, {casted, stream});
 
   rewriter.replaceOp(deallocOp, {stream});
   return success();
 }
 
 static bool isGpuAsyncTokenType(Value value) {
   return value.getType().isa<gpu::AsyncTokenType>();
 }
 
 // Converts !gpu.async.token operands of `async.yield` to runtime calls. The
 // !gpu.async.token are lowered to stream within the async.execute region, but
 // are passed as events between them. For each !gpu.async.token operand, we
 // create an event and record it on the stream.
 LogicalResult ConvertAsyncYieldToGpuRuntimeCallPattern::matchAndRewrite(
-    async::YieldOp yieldOp, ArrayRef<Value> operands,
+    async::YieldOp yieldOp, OpAdaptor adaptor,
     ConversionPatternRewriter &rewriter) const {
   if (llvm::none_of(yieldOp.operands(), isGpuAsyncTokenType))
     return rewriter.notifyMatchFailure(yieldOp, "no gpu async token operand");
 
   Location loc = yieldOp.getLoc();
-  SmallVector<Value, 4> newOperands(operands.begin(), operands.end());
+  SmallVector<Value, 4> newOperands(adaptor.getOperands());
   llvm::SmallDenseSet<Value> streams;
   for (auto &operand : yieldOp->getOpOperands()) {
     if (!isGpuAsyncTokenType(operand.get()))
       continue;
     auto idx = operand.getOperandNumber();
-    auto stream = operands[idx];
+    auto stream = adaptor.getOperands()[idx];
     auto event = eventCreateCallBuilder.create(loc, rewriter, {}).getResult(0);
     eventRecordCallBuilder.create(loc, rewriter, {event, stream});
     newOperands[idx] = event;
     streams.insert(stream);
   }
   for (auto stream : streams)
     streamDestroyCallBuilder.create(loc, rewriter, {stream});
 
static bool isDefinedByCallTo(Value value, StringRef functionName) {
   return false;
 }
 
 // Converts `gpu.wait` to runtime calls. The converted op synchronizes the host
 // with the stream/event operands. The operands are destroyed. That is, it
 // assumes that it is not used afterwards or elsewhere. Otherwise we will get a
 // runtime error. Eventually, we should guarantee this property.
 LogicalResult ConvertWaitOpToGpuRuntimeCallPattern::matchAndRewrite(
-    gpu::WaitOp waitOp, ArrayRef<Value> operands,
+    gpu::WaitOp waitOp, OpAdaptor adaptor,
     ConversionPatternRewriter &rewriter) const {
   if (waitOp.asyncToken())
     return rewriter.notifyMatchFailure(waitOp, "Cannot convert async op.");
 
   Location loc = waitOp.getLoc();
 
-  for (auto operand : operands) {
+  for (auto operand : adaptor.getOperands()) {
     if (isDefinedByCallTo(operand, streamCreateCallBuilder.functionName)) {
       // The converted operand's definition created a stream.
       streamSynchronizeCallBuilder.create(loc, rewriter, {operand});
       streamDestroyCallBuilder.create(loc, rewriter, {operand});
     } else {
       // Otherwise the converted operand is an event. This assumes that we use
       // events in control flow code as well.
       eventSynchronizeCallBuilder.create(loc, rewriter, {operand});
       eventDestroyCallBuilder.create(loc, rewriter, {operand});
     }
   }
 
   rewriter.eraseOp(waitOp);
   return success();
 }
 
 // Converts `gpu.wait async` to runtime calls. The converted op creates a new
 // stream that is synchronized with stream/event operands. The operands are
 // destroyed. That is, it assumes that it is not used afterwards or elsewhere.
 // Otherwise we will get a runtime error. Eventually, we should guarantee this
 // property.
 LogicalResult ConvertWaitAsyncOpToGpuRuntimeCallPattern::matchAndRewrite(
-    gpu::WaitOp waitOp, ArrayRef<Value> operands,
+    gpu::WaitOp waitOp, OpAdaptor adaptor,
     ConversionPatternRewriter &rewriter) const {
   if (!waitOp.asyncToken())
     return rewriter.notifyMatchFailure(waitOp, "Can only convert async op.");
 
   Location loc = waitOp.getLoc();
 
   auto insertionPoint = rewriter.saveInsertionPoint();
   SmallVector<Value, 1> events;
-  for (auto pair : llvm::zip(waitOp.asyncDependencies(), operands)) {
+  for (auto pair :
+       llvm::zip(waitOp.asyncDependencies(), adaptor.getOperands())) {
     auto operand = std::get<1>(pair);
     if (isDefinedByCallTo(operand, streamCreateCallBuilder.functionName)) {
       // The converted operand's definition created a stream. Insert an event
       // into the stream just after the last use of the original token operand.
       auto *defOp = std::get<0>(pair).getDefiningOp();
       rewriter.setInsertionPointAfter(defOp);
       auto event =
           eventCreateCallBuilder.create(loc, rewriter, {}).getResult(0);

 // %array = alloca(NumParameters * sizeof(void *))
 // for (i : [0, NumParameters))
 //   %fieldPtr = llvm.getelementptr %struct[0, i]
 //   llvm.store parameters[i], %fieldPtr
 //   %elementPtr = llvm.getelementptr %array[i]
 //   llvm.store %fieldPtr, %elementPtr
 // return %array
 Value ConvertLaunchFuncOpToGpuRuntimeCallPattern::generateParamsArray(
-    gpu::LaunchFuncOp launchOp, ArrayRef<Value> operands,
-    OpBuilder &builder) const {
+    gpu::LaunchFuncOp launchOp, OpAdaptor adaptor, OpBuilder &builder) const {
   auto loc = launchOp.getLoc();
   auto numKernelOperands = launchOp.getNumKernelOperands();
   auto arguments = getTypeConverter()->promoteOperands(
       loc, launchOp.getOperands().take_back(numKernelOperands),
-      operands.take_back(numKernelOperands), builder);
+      adaptor.getOperands().take_back(numKernelOperands), builder);
   auto numArguments = arguments.size();
   SmallVector<Type, 4> argumentTypes;
   argumentTypes.reserve(numArguments);
   for (auto argument : arguments)
     argumentTypes.push_back(argument.getType());
   auto structType = LLVM::LLVMStructType::getNewIdentified(context, StringRef(),
                                                            argumentTypes);
   auto one = builder.create<LLVM::ConstantOp>(loc, llvmInt32Type,

 // call %launchKernel(%3, <launchOp operands 0..5>, 0, %4, %5, nullptr)
 // call %streamSynchronize(%4)
 // call %streamDestroy(%4)
 // call %moduleUnload(%1)
 //
 // If the op is async, the stream corresponds to the (single) async dependency
 // as well as the async token the op produces.
 LogicalResult ConvertLaunchFuncOpToGpuRuntimeCallPattern::matchAndRewrite(
-    gpu::LaunchFuncOp launchOp, ArrayRef<Value> operands,
+    gpu::LaunchFuncOp launchOp, OpAdaptor adaptor,
     ConversionPatternRewriter &rewriter) const {
-  if (failed(areAllLLVMTypes(launchOp, operands, rewriter)))
+  if (failed(areAllLLVMTypes(launchOp, adaptor.getOperands(), rewriter)))
     return failure();
 
   if (launchOp.asyncDependencies().size() > 1)
     return rewriter.notifyMatchFailure(
         launchOp, "Cannot convert with more than one async dependency.");
 
   // Fail when the synchronous version of the op has async dependencies. The
   // lowering destroys the stream, and we do not want to check that there is no
LogicalResult ConvertLaunchFuncOpToGpuRuntimeCallPattern::matchAndRewrite(
   // the kernel function.
   auto kernelName = generateKernelNameConstant(
       launchOp.getKernelModuleName().getValue(),
       launchOp.getKernelName().getValue(), loc, rewriter);
   auto function = moduleGetFunctionCallBuilder.create(
       loc, rewriter, {module.getResult(0), kernelName});
   auto zero = rewriter.create<LLVM::ConstantOp>(loc, llvmInt32Type,
                                                 rewriter.getI32IntegerAttr(0));
-  auto adaptor =
-      gpu::LaunchFuncOpAdaptor(operands, launchOp->getAttrDictionary());
   Value stream =
       adaptor.asyncDependencies().empty()
           ? streamCreateCallBuilder.create(loc, rewriter, {}).getResult(0)
           : adaptor.asyncDependencies().front();
   // Create array of pointers to kernel arguments.
-  auto kernelParams = generateParamsArray(launchOp, operands, rewriter);
+  auto kernelParams = generateParamsArray(launchOp, adaptor, rewriter);
   auto nullpointer = rewriter.create<LLVM::NullOp>(loc, llvmPointerPointerType);
   launchKernelCallBuilder.create(loc, rewriter,
                                  {function.getResult(0), adaptor.gridSizeX(),
                                   adaptor.gridSizeY(), adaptor.gridSizeZ(),
                                   adaptor.blockSizeX(), adaptor.blockSizeY(),
                                   adaptor.blockSizeZ(),
                                   /*sharedMemBytes=*/zero, stream, kernelParams,
                                   /*extra=*/nullpointer});
if (launchOp.asyncToken()) {
     rewriter.eraseOp(launchOp);
   }
   moduleUnloadCallBuilder.create(loc, rewriter, module.getResult(0));
 
   return success();
 }
 
 LogicalResult ConvertMemcpyOpToGpuRuntimeCallPattern::matchAndRewrite(
-    gpu::MemcpyOp memcpyOp, ArrayRef<Value> operands,
+    gpu::MemcpyOp memcpyOp, OpAdaptor adaptor,
     ConversionPatternRewriter &rewriter) const {
   auto memRefType = memcpyOp.src().getType().cast<MemRefType>();
 
-  if (failed(areAllLLVMTypes(memcpyOp, operands, rewriter)) ||
+  if (failed(areAllLLVMTypes(memcpyOp, adaptor.getOperands(), rewriter)) ||
       !isConvertibleAndHasIdentityMaps(memRefType) ||
       failed(isAsyncWithOneDependency(rewriter, memcpyOp)))
     return failure();
 
   auto loc = memcpyOp.getLoc();
-  auto adaptor = gpu::MemcpyOpAdaptor(operands, memcpyOp->getAttrDictionary());
 
   MemRefDescriptor srcDesc(adaptor.src());
   Value numElements = getNumElements(rewriter, loc, memRefType, srcDesc);
 
   Type elementPtrType = getElementPtrType(memRefType);
   Value nullPtr = rewriter.create<LLVM::NullOp>(loc, elementPtrType);
   Value gepPtr = rewriter.create<LLVM::GEPOp>(
       loc, elementPtrType, ArrayRef<Value>{nullPtr, numElements});
LogicalResult ConvertMemcpyOpToGpuRuntimeCallPattern::matchAndRewrite(
   memcpyCallBuilder.create(loc, rewriter, {dst, src, sizeBytes, stream});
 
   rewriter.replaceOp(memcpyOp, {stream});
 
   return success();
 }
 
 LogicalResult ConvertMemsetOpToGpuRuntimeCallPattern::matchAndRewrite(
-    gpu::MemsetOp memsetOp, ArrayRef<Value> operands,
+    gpu::MemsetOp memsetOp, OpAdaptor adaptor,
     ConversionPatternRewriter &rewriter) const {
   auto memRefType = memsetOp.dst().getType().cast<MemRefType>();
 
-  if (failed(areAllLLVMTypes(memsetOp, operands, rewriter)) ||
+  if (failed(areAllLLVMTypes(memsetOp, adaptor.getOperands(), rewriter)) ||
       !isConvertibleAndHasIdentityMaps(memRefType) ||
       failed(isAsyncWithOneDependency(rewriter, memsetOp)))
     return failure();
 
   auto loc = memsetOp.getLoc();
-  auto adaptor = gpu::MemsetOpAdaptor(operands, memsetOp->getAttrDictionary());
 
   Type valueType = adaptor.value().getType();
   if (!valueType.isIntOrFloat() || valueType.getIntOrFloatBitWidth() != 32) {
     return rewriter.notifyMatchFailure(memsetOp,
                                        "value must be a 32 bit scalar");
   }
 
   MemRefDescriptor dstDesc(adaptor.dst());

mlir/lib/Conversion/GPUCommon/IndexIntrinsicsOpLowering.h

 
 public:
   explicit GPUIndexIntrinsicOpLowering(LLVMTypeConverter &typeConverter)
       : ConvertOpToLLVMPattern<Op>(typeConverter),
         indexBitwidth(typeConverter.getIndexTypeBitwidth()) {}
 
   // Convert the kernel arguments to an LLVM type, preserve the rest.
   LogicalResult
-  matchAndRewrite(Op op, ArrayRef<Value> operands,
+  matchAndRewrite(Op op, typename Op::Adaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     auto loc = op->getLoc();
     MLIRContext *context = rewriter.getContext();
     Value newOp;
     switch (dimensionToIndex(op)) {
     case X:
       newOp = rewriter.create<XOp>(loc, IntegerType::get(context, 32));
       break;

mlir/lib/Conversion/GPUCommon/OpToFuncCallLowering.h

 struct OpToFuncCallLowering : public ConvertOpToLLVMPattern<SourceOp> {
 public:
   explicit OpToFuncCallLowering(LLVMTypeConverter &lowering_, StringRef f32Func,
                                 StringRef f64Func)
       : ConvertOpToLLVMPattern<SourceOp>(lowering_), f32Func(f32Func),
         f64Func(f64Func) {}
 
   LogicalResult
-  matchAndRewrite(SourceOp op, ArrayRef<Value> operands,
+  matchAndRewrite(SourceOp op, typename SourceOp::Adaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     using LLVM::LLVMFuncOp;
 
     static_assert(
         std::is_base_of<OpTrait::OneResult<SourceOp>, SourceOp>::value,
         "expected single result op");
 
     static_assert(std::is_base_of<OpTrait::SameOperandsAndResultType<SourceOp>,
                                   SourceOp>::value,
                   "expected op with same operand and result types");
 
     SmallVector<Value, 1> castedOperands;
-    for (Value operand : operands)
+    for (Value operand : adaptor.getOperands())
       castedOperands.push_back(maybeCast(operand, rewriter));
 
     Type resultType = castedOperands.front().getType();
     Type funcType = getFunctionType(resultType, castedOperands);
     StringRef funcName = getFunctionName(
         funcType.cast<LLVM::LLVMFunctionType>().getReturnType());
     if (funcName.empty())
       return failure();
 
     LLVMFuncOp funcOp = appendOrGetFuncOp(funcName, funcType, op);
     auto callOp = rewriter.create<LLVM::CallOp>(
         op->getLoc(), resultType, SymbolRefAttr::get(funcOp), castedOperands);
 
-    if (resultType == operands.front().getType()) {
+    if (resultType == adaptor.getOperands().front().getType()) {
       rewriter.replaceOp(op, {callOp.getResult(0)});
       return success();
     }
 
     Value truncated = rewriter.create<LLVM::FPTruncOp>(
-        op->getLoc(), operands.front().getType(), callOp.getResult(0));
+        op->getLoc(), adaptor.getOperands().front().getType(),
+        callOp.getResult(0));
     rewriter.replaceOp(op, {truncated});
     return success();
   }
 
 private:
   Value maybeCast(Value operand, PatternRewriter &rewriter) const {
     Type type = operand.getType();
     if (!type.isa<Float16Type>())
       return operand;
 
     return rewriter.create<LLVM::FPExtOp>(
         operand.getLoc(), Float32Type::get(rewriter.getContext()), operand);
   }
 
-  Type getFunctionType(Type resultType, ArrayRef<Value> operands) const {
-    SmallVector<Type, 1> operandTypes;
-    for (Value operand : operands) {
-      operandTypes.push_back(operand.getType());
-    }
+  Type getFunctionType(Type resultType, ValueRange operands) const {
+    SmallVector<Type> operandTypes(operands.getTypes());
     return LLVM::LLVMFunctionType::get(resultType, operandTypes);
   }
 
   StringRef getFunctionName(Type type) const {
     if (type.isa<Float32Type>())
       return f32Func;
     if (type.isa<Float64Type>())
       return f64Func;

mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp

struct GPUShuffleOpLowering : public ConvertOpToLLVMPattern<gpu::ShuffleOp> {
   ///     %mask_and_clamp = llvm.sub %width, %one : 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(gpu::ShuffleOp op, ArrayRef<Value> operands,
+  matchAndRewrite(gpu::ShuffleOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     Location loc = op->getLoc();
-    gpu::ShuffleOpAdaptor adaptor(operands);
 
     auto valueTy = adaptor.value().getType();
     auto int32Type = IntegerType::get(rewriter.getContext(), 32);
     auto predTy = IntegerType::get(rewriter.getContext(), 1);
     auto resultTy = LLVM::LLVMStructType::getLiteral(rewriter.getContext(),
                                                      {valueTy, predTy});
 
     Value one = rewriter.create<LLVM::ConstantOp>(

mlir/lib/Conversion/GPUToNVVM/WmmaOpsToNvvm.cpp

 /// after it has been loaded.
 struct WmmaLoadOpToNVVMLowering
     : public ConvertOpToLLVMPattern<gpu::SubgroupMmaLoadMatrixOp> {
   using ConvertOpToLLVMPattern<
       gpu::SubgroupMmaLoadMatrixOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
   matchAndRewrite(gpu::SubgroupMmaLoadMatrixOp subgroupMmaLoadMatrixOp,
-                  ArrayRef<Value> operands,
+                  OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     Operation *op = subgroupMmaLoadMatrixOp.getOperation();
-    if (failed(areAllLLVMTypes(op, operands, rewriter)))
+    if (failed(areAllLLVMTypes(op, adaptor.getOperands(), rewriter)))
       return failure();
 
     unsigned indexTypeBitwidth =
         this->getTypeConverter()->getIndexTypeBitwidth();
 
     // The corresponding intrinsics expects leadDimension to be a 32-bit
     // integer, so all the calculations of linearizing the load address
     // must also follow this restriction.
     if (indexTypeBitwidth != 32)
       return rewriter.notifyMatchFailure(
           op, "Expected indices to the memref to be 32-bit wide.");
     Location loc = op->getLoc();
 
     auto leadDimension = subgroupMmaLoadMatrixOp.leadDimensionAttr();
 
-    gpu::SubgroupMmaLoadMatrixOpAdaptor adaptor(operands);
     // MemRefDescriptor to extract alignedPtr and offset.
     MemRefDescriptor promotedSrcOp(adaptor.srcMemref());
 
     // Emit ops which compute the load offset using `srcOffsetI`,
     // `srcOffsetJ`. The actualOffset is (memrefOffset + (alignedPtr +
     // ((leadDimension * srcOffsetI) + srcOffsetJ)). The memrefs here are
     // assumed to be normalized and hence the simple conversion works.
     SmallVector<Value> indices(adaptor.indices());

 /// convert the data in the format that is needed by the NVVM op.
 struct WmmaStoreOpToNVVMLowering
     : public ConvertOpToLLVMPattern<gpu::SubgroupMmaStoreMatrixOp> {
   using ConvertOpToLLVMPattern<
       gpu::SubgroupMmaStoreMatrixOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
   matchAndRewrite(gpu::SubgroupMmaStoreMatrixOp subgroupMmaStoreMatrixOp,
-                  ArrayRef<Value> operands,
+                  OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     Operation *op = subgroupMmaStoreMatrixOp.getOperation();
-    if (failed(areAllLLVMTypes(op, operands, rewriter)))
+    if (failed(areAllLLVMTypes(op, adaptor.getOperands(), rewriter)))
       return failure();
 
     unsigned indexTypeBitwidth =
         this->getTypeConverter()->getIndexTypeBitwidth();
     // The corresponding intrinsics expects leadDimension to be a 32-bit
     // integer, so all the calculations of linearizing the store address
     // must also follow this restriction.
     if (indexTypeBitwidth != 32)
       return rewriter.notifyMatchFailure(
           op, "expected indices to the memref to be 32-bit wide.");
 
     Location loc = op->getLoc();
 
-    gpu::SubgroupMmaStoreMatrixOpAdaptor adaptor(operands);
     // MemRefDescriptor to extract alignedPtr and offset.
     MemRefDescriptor promotedDstOp(adaptor.dstMemref());
 
     auto leadDimension = subgroupMmaStoreMatrixOp.leadDimensionAttr();
 
     // Emit ops which compute the store offset using `dstOffsetI`,
     // `dstOffsetJ`. The actualOffset is (memrefOffset + (alignedPtr +
     // ((leadDimension * dstOffsetI) + dstOffsetJ)).

 /// in the NVVM dialect.
 struct WmmaMmaOpToNVVMLowering
     : public ConvertOpToLLVMPattern<gpu::SubgroupMmaComputeOp> {
   using ConvertOpToLLVMPattern<
       gpu::SubgroupMmaComputeOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
   matchAndRewrite(gpu::SubgroupMmaComputeOp subgroupMmaComputeOp,
-                  ArrayRef<Value> operands,
+                  OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     Operation *op = subgroupMmaComputeOp.getOperation();
-    if (failed(areAllLLVMTypes(op, operands, rewriter)))
+    if (failed(areAllLLVMTypes(op, adaptor.getOperands(), rewriter)))
       return failure();
 
     Location loc = op->getLoc();
 
     // The wmma.mma intrinsic in llvm requires the operands as individual
     // values. So individual elements from the memrefs need to be extracted and
     // then passed on to the intrinsic call. Emit llvm ops to extract individual
     // values form lowered memrefs.
matchAndRewrite(gpu::SubgroupMmaComputeOp subgroupMmaComputeOp,
     ArrayRef<int64_t> aTypeShape = aType.getShape();
     gpu::MMAMatrixType bType =
         subgroupMmaComputeOp.opB().getType().cast<gpu::MMAMatrixType>();
     ArrayRef<int64_t> bTypeShape = bType.getShape();
     gpu::MMAMatrixType cType =
         subgroupMmaComputeOp.opC().getType().cast<gpu::MMAMatrixType>();
     ArrayRef<int64_t> cTypeShape = cType.getShape();
 
-    gpu::SubgroupMmaComputeOpAdaptor transformedOperands(operands);
-    unpackOp(transformedOperands.opA());
-    unpackOp(transformedOperands.opB());
-    unpackOp(transformedOperands.opC());
+    unpackOp(adaptor.opA());
+    unpackOp(adaptor.opB());
+    unpackOp(adaptor.opC());
 
     if (cType.getElementType().isF16()) {
       if (aTypeShape[0] == 16 && aTypeShape[1] == 16 && bTypeShape[0] == 16 &&
           bTypeShape[1] == 16 && cTypeShape[0] == 16 && cTypeShape[1] == 16) {
         // Create nvvm.wmma.mma op.
         rewriter.replaceOpWithNewOp<NVVM::WMMAMmaF16F16M16N16K16Op>(
-            op, transformedOperands.opC().getType(), unpackedOps);
+            op, adaptor.opC().getType(), unpackedOps);
 
         return success();
       }
       return rewriter.notifyMatchFailure(op, kInvalidCaseStr);
     }
     if (cType.getElementType().isF32()) {
       if (aTypeShape[0] == 16 && aTypeShape[1] == 16 && bTypeShape[0] == 16 &&
           bTypeShape[1] == 16 && cTypeShape[0] == 16 && cTypeShape[1] == 16) {
         // Create nvvm.wmma.mma op.
         rewriter.replaceOpWithNewOp<NVVM::WMMAMmaF32F32M16N16K16Op>(
-            op, transformedOperands.opC().getType(), unpackedOps);
+            op, adaptor.opC().getType(), unpackedOps);
 
         return success();
       }
       return rewriter.notifyMatchFailure(op, kInvalidCaseStr);
     }
     return failure();
   }
 };
 
 /// Convert GPU MMA ConstantMatrixOp to a chain of InsertValueOp.
 struct WmmaConstantOpToNVVMLowering
     : public ConvertOpToLLVMPattern<gpu::SubgroupMmaConstantMatrixOp> {
   using ConvertOpToLLVMPattern<
       gpu::SubgroupMmaConstantMatrixOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
   matchAndRewrite(gpu::SubgroupMmaConstantMatrixOp subgroupMmaConstantOp,
-                  ArrayRef<Value> operands,
+                  OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    if (failed(areAllLLVMTypes(subgroupMmaConstantOp.getOperation(), operands,
-                               rewriter)))
+    if (failed(areAllLLVMTypes(subgroupMmaConstantOp.getOperation(),
+                               adaptor.getOperands(), rewriter)))
       return failure();
     Location loc = subgroupMmaConstantOp.getLoc();
-    Value cst = operands[0];
+    Value cst = adaptor.getOperands()[0];
     LLVM::LLVMStructType type = convertMMAToLLVMType(
         subgroupMmaConstantOp.getType().cast<gpu::MMAMatrixType>());
     // If the element type is a vector create a vector from the operand.
     if (auto vecType = type.getBody()[0].dyn_cast<VectorType>()) {
       Value vecCst = rewriter.create<LLVM::UndefOp>(loc, vecType);
       for (int64_t vecEl = 0; vecEl < vecType.getNumElements(); vecEl++) {
         Value idx = rewriter.create<LLVM::ConstantOp>(
             loc, typeConverter->convertType(rewriter.getIntegerType(32)),

mlir/lib/Conversion/LinalgToLLVM/LinalgToLLVM.cpp

 
 namespace {
 // RangeOp creates a new range descriptor.
 class RangeOpConversion : public ConvertOpToLLVMPattern<RangeOp> {
 public:
   using ConvertOpToLLVMPattern<RangeOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(RangeOp rangeOp, ArrayRef<Value> operands,
+  matchAndRewrite(RangeOp rangeOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     auto rangeDescriptorTy = convertRangeType(
         rangeOp.getType().cast<RangeType>(), *getTypeConverter());
 
     ImplicitLocOpBuilder b(rangeOp->getLoc(), rewriter);
 
     // Fill in an aggregate value of the descriptor.
-    RangeOpAdaptor adaptor(operands);
     Value desc = b.create<LLVM::UndefOp>(rangeDescriptorTy);
     desc = b.create<LLVM::InsertValueOp>(desc, adaptor.min(),
                                          rewriter.getI64ArrayAttr(0));
     desc = b.create<LLVM::InsertValueOp>(desc, adaptor.max(),
                                          rewriter.getI64ArrayAttr(1));
     desc = b.create<LLVM::InsertValueOp>(desc, adaptor.step(),
                                          rewriter.getI64ArrayAttr(2));
     rewriter.replaceOp(rangeOp, desc);
     return success();
   }
 };
 
 
 // YieldOp produces and LLVM::ReturnOp.
 class YieldOpConversion : public ConvertOpToLLVMPattern<linalg::YieldOp> {
 public:
   using ConvertOpToLLVMPattern<linalg::YieldOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(linalg::YieldOp op, ArrayRef<Value> operands,
+  matchAndRewrite(linalg::YieldOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    rewriter.replaceOpWithNewOp<LLVM::ReturnOp>(op, operands);
+    rewriter.replaceOpWithNewOp<LLVM::ReturnOp>(op, adaptor.getOperands());
     return success();
   }
 };
 } // namespace
 
 /// Populate the given list with patterns that convert from Linalg to LLVM.
 void mlir::populateLinalgToLLVMConversionPatterns(LLVMTypeConverter &converter,
                                                   RewritePatternSet &patterns) {

mlir/lib/Conversion/MathToLLVM/MathToLLVM.cpp

 using SinOpLowering = VectorConvertToLLVMPattern<math::SinOp, LLVM::SinOp>;
 using SqrtOpLowering = VectorConvertToLLVMPattern<math::SqrtOp, LLVM::SqrtOp>;
 
 // A `expm1` is converted into `exp - 1`.
 struct ExpM1OpLowering : public ConvertOpToLLVMPattern<math::ExpM1Op> {
   using ConvertOpToLLVMPattern<math::ExpM1Op>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(math::ExpM1Op op, ArrayRef<Value> operands,
+  matchAndRewrite(math::ExpM1Op op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    math::ExpM1Op::Adaptor transformed(operands);
-    auto operandType = transformed.operand().getType();
+    auto operandType = adaptor.operand().getType();
 
     if (!operandType || !LLVM::isCompatibleType(operandType))
       return failure();
 
     auto loc = op.getLoc();
     auto resultType = op.getResult().getType();
     auto floatType = getElementTypeOrSelf(resultType).cast<FloatType>();
     auto floatOne = rewriter.getFloatAttr(floatType, 1.0);
 
     if (!operandType.isa<LLVM::LLVMArrayType>()) {
       LLVM::ConstantOp one;
       if (LLVM::isCompatibleVectorType(operandType)) {
         one = rewriter.create<LLVM::ConstantOp>(
             loc, operandType,
             SplatElementsAttr::get(resultType.cast<ShapedType>(), floatOne));
       } else {
         one = rewriter.create<LLVM::ConstantOp>(loc, operandType, floatOne);
       }
-      auto exp = rewriter.create<LLVM::ExpOp>(loc, transformed.operand());
+      auto exp = rewriter.create<LLVM::ExpOp>(loc, adaptor.operand());
       rewriter.replaceOpWithNewOp<LLVM::FSubOp>(op, operandType, exp, one);
       return success();
     }
 
     auto vectorType = resultType.dyn_cast<VectorType>();
     if (!vectorType)
       return rewriter.notifyMatchFailure(op, "expected vector result type");
 
     return LLVM::detail::handleMultidimensionalVectors(
-        op.getOperation(), operands, *getTypeConverter(),
+        op.getOperation(), adaptor.getOperands(), *getTypeConverter(),
         [&](Type llvm1DVectorTy, ValueRange operands) {
           auto splatAttr = SplatElementsAttr::get(
               mlir::VectorType::get(
                   {LLVM::getVectorNumElements(llvm1DVectorTy).getFixedValue()},
                   floatType),
               floatOne);
           auto one =
               rewriter.create<LLVM::ConstantOp>(loc, llvm1DVectorTy, splatAttr);
           auto exp =
               rewriter.create<LLVM::ExpOp>(loc, llvm1DVectorTy, operands[0]);
           return rewriter.create<LLVM::FSubOp>(loc, llvm1DVectorTy, exp, one);
         },
         rewriter);
   }
 };
 
 // A `log1p` is converted into `log(1 + ...)`.
 struct Log1pOpLowering : public ConvertOpToLLVMPattern<math::Log1pOp> {
   using ConvertOpToLLVMPattern<math::Log1pOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(math::Log1pOp op, ArrayRef<Value> operands,
+  matchAndRewrite(math::Log1pOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    math::Log1pOp::Adaptor transformed(operands);
-    auto operandType = transformed.operand().getType();
+    auto operandType = adaptor.operand().getType();
 
     if (!operandType || !LLVM::isCompatibleType(operandType))
       return rewriter.notifyMatchFailure(op, "unsupported operand type");
 
     auto loc = op.getLoc();
     auto resultType = op.getResult().getType();
     auto floatType = getElementTypeOrSelf(resultType).cast<FloatType>();
     auto floatOne = rewriter.getFloatAttr(floatType, 1.0);
 
     if (!operandType.isa<LLVM::LLVMArrayType>()) {
       LLVM::ConstantOp one =
           LLVM::isCompatibleVectorType(operandType)
               ? rewriter.create<LLVM::ConstantOp>(
                     loc, operandType,
                     SplatElementsAttr::get(resultType.cast<ShapedType>(),
                                            floatOne))
               : rewriter.create<LLVM::ConstantOp>(loc, operandType, floatOne);
 
       auto add = rewriter.create<LLVM::FAddOp>(loc, operandType, one,
-                                               transformed.operand());
+                                               adaptor.operand());
       rewriter.replaceOpWithNewOp<LLVM::LogOp>(op, operandType, add);
       return success();
     }
 
     auto vectorType = resultType.dyn_cast<VectorType>();
     if (!vectorType)
       return rewriter.notifyMatchFailure(op, "expected vector result type");
 
     return LLVM::detail::handleMultidimensionalVectors(
-        op.getOperation(), operands, *getTypeConverter(),
+        op.getOperation(), adaptor.getOperands(), *getTypeConverter(),
         [&](Type llvm1DVectorTy, ValueRange operands) {
           auto splatAttr = SplatElementsAttr::get(
               mlir::VectorType::get(
                   {LLVM::getVectorNumElements(llvm1DVectorTy).getFixedValue()},
                   floatType),
               floatOne);
           auto one =
               rewriter.create<LLVM::ConstantOp>(loc, llvm1DVectorTy, splatAttr);
           auto add = rewriter.create<LLVM::FAddOp>(loc, llvm1DVectorTy, one,
                                                    operands[0]);
           return rewriter.create<LLVM::LogOp>(loc, llvm1DVectorTy, add);
         },
         rewriter);
   }
 };
 
 // A `rsqrt` is converted into `1 / sqrt`.
 struct RsqrtOpLowering : public ConvertOpToLLVMPattern<math::RsqrtOp> {
   using ConvertOpToLLVMPattern<math::RsqrtOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(math::RsqrtOp op, ArrayRef<Value> operands,
+  matchAndRewrite(math::RsqrtOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    math::RsqrtOp::Adaptor transformed(operands);
-    auto operandType = transformed.operand().getType();
+    auto operandType = adaptor.operand().getType();
 
     if (!operandType || !LLVM::isCompatibleType(operandType))
       return failure();
 
     auto loc = op.getLoc();
     auto resultType = op.getResult().getType();
     auto floatType = getElementTypeOrSelf(resultType).cast<FloatType>();
     auto floatOne = rewriter.getFloatAttr(floatType, 1.0);
 
     if (!operandType.isa<LLVM::LLVMArrayType>()) {
       LLVM::ConstantOp one;
       if (LLVM::isCompatibleVectorType(operandType)) {
         one = rewriter.create<LLVM::ConstantOp>(
             loc, operandType,
             SplatElementsAttr::get(resultType.cast<ShapedType>(), floatOne));
       } else {
         one = rewriter.create<LLVM::ConstantOp>(loc, operandType, floatOne);
       }
-      auto sqrt = rewriter.create<LLVM::SqrtOp>(loc, transformed.operand());
+      auto sqrt = rewriter.create<LLVM::SqrtOp>(loc, adaptor.operand());
       rewriter.replaceOpWithNewOp<LLVM::FDivOp>(op, operandType, one, sqrt);
       return success();
     }
 
     auto vectorType = resultType.dyn_cast<VectorType>();
     if (!vectorType)
       return failure();
 
     return LLVM::detail::handleMultidimensionalVectors(
-        op.getOperation(), operands, *getTypeConverter(),
+        op.getOperation(), adaptor.getOperands(), *getTypeConverter(),
         [&](Type llvm1DVectorTy, ValueRange operands) {
           auto splatAttr = SplatElementsAttr::get(
               mlir::VectorType::get(
                   {LLVM::getVectorNumElements(llvm1DVectorTy).getFixedValue()},
                   floatType),
               floatOne);
           auto one =
               rewriter.create<LLVM::ConstantOp>(loc, llvm1DVectorTy, splatAttr);

mlir/lib/Conversion/MemRefToLLVM/MemRefToLLVM.cpp

struct AllocaOpLowering : public AllocLikeOpLLVMLowering {
   }
 };
 
 struct AllocaScopeOpLowering
     : public ConvertOpToLLVMPattern<memref::AllocaScopeOp> {
   using ConvertOpToLLVMPattern<memref::AllocaScopeOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(memref::AllocaScopeOp allocaScopeOp, ArrayRef<Value> operands,
+  matchAndRewrite(memref::AllocaScopeOp allocaScopeOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     OpBuilder::InsertionGuard guard(rewriter);
     Location loc = allocaScopeOp.getLoc();
 
     // Split the current block before the AllocaScopeOp to create the inlining
     // point.
     auto *currentBlock = rewriter.getInsertionBlock();
     auto *remainingOpsBlock =

 };
 
 struct AssumeAlignmentOpLowering
     : public ConvertOpToLLVMPattern<memref::AssumeAlignmentOp> {
   using ConvertOpToLLVMPattern<
       memref::AssumeAlignmentOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(memref::AssumeAlignmentOp op, ArrayRef<Value> operands,
+  matchAndRewrite(memref::AssumeAlignmentOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    memref::AssumeAlignmentOp::Adaptor transformed(operands);
-    Value memref = transformed.memref();
+    Value memref = adaptor.memref();
     unsigned alignment = op.alignment();
     auto loc = op.getLoc();
 
     MemRefDescriptor memRefDescriptor(memref);
     Value ptr = memRefDescriptor.alignedPtr(rewriter, memref.getLoc());
 
     // Emit llvm.assume(memref.alignedPtr & (alignment - 1) == 0). Notice that
     // the asserted memref.alignedPtr isn't used anywhere else, as the real

 // in any way.
 struct DeallocOpLowering : public ConvertOpToLLVMPattern<memref::DeallocOp> {
   using ConvertOpToLLVMPattern<memref::DeallocOp>::ConvertOpToLLVMPattern;
 
   explicit DeallocOpLowering(LLVMTypeConverter &converter)
       : ConvertOpToLLVMPattern<memref::DeallocOp>(converter) {}
 
   LogicalResult
-  matchAndRewrite(memref::DeallocOp op, ArrayRef<Value> operands,
+  matchAndRewrite(memref::DeallocOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    assert(operands.size() == 1 && "dealloc takes one operand");
-    memref::DeallocOp::Adaptor transformed(operands);
-
     // Insert the `free` declaration if it is not already present.
     auto freeFunc = LLVM::lookupOrCreateFreeFn(op->getParentOfType<ModuleOp>());
-    MemRefDescriptor memref(transformed.memref());
+    MemRefDescriptor memref(adaptor.memref());
     Value casted = rewriter.create<LLVM::BitcastOp>(
         op.getLoc(), getVoidPtrType(),
         memref.allocatedPtr(rewriter, op.getLoc()));
     rewriter.replaceOpWithNewOp<LLVM::CallOp>(
         op, TypeRange(), SymbolRefAttr::get(freeFunc), casted);
     return success();
   }
 };
 
 // A `dim` is converted to a constant for static sizes and to an access to the
 // size stored in the memref descriptor for dynamic sizes.
 struct DimOpLowering : public ConvertOpToLLVMPattern<memref::DimOp> {
   using ConvertOpToLLVMPattern<memref::DimOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(memref::DimOp dimOp, ArrayRef<Value> operands,
+  matchAndRewrite(memref::DimOp dimOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     Type operandType = dimOp.source().getType();
     if (operandType.isa<UnrankedMemRefType>()) {
-      rewriter.replaceOp(dimOp, {extractSizeOfUnrankedMemRef(
-                                    operandType, dimOp, operands, rewriter)});
+      rewriter.replaceOp(
+          dimOp, {extractSizeOfUnrankedMemRef(
+                     operandType, dimOp, adaptor.getOperands(), rewriter)});
 
       return success();
     }
     if (operandType.isa<MemRefType>()) {
-      rewriter.replaceOp(dimOp, {extractSizeOfRankedMemRef(
-                                    operandType, dimOp, operands, rewriter)});
+      rewriter.replaceOp(
+          dimOp, {extractSizeOfRankedMemRef(operandType, dimOp,
+                                            adaptor.getOperands(), rewriter)});
       return success();
     }
     llvm_unreachable("expected MemRefType or UnrankedMemRefType");
   }
 
 private:
   Value extractSizeOfUnrankedMemRef(Type operandType, memref::DimOp dimOp,
-                                    ArrayRef<Value> operands,
+                                    OpAdaptor adaptor,
                                     ConversionPatternRewriter &rewriter) const {
     Location loc = dimOp.getLoc();
-    memref::DimOp::Adaptor transformed(operands);
 
     auto unrankedMemRefType = operandType.cast<UnrankedMemRefType>();
     auto scalarMemRefType =
         MemRefType::get({}, unrankedMemRefType.getElementType());
     unsigned addressSpace = unrankedMemRefType.getMemorySpaceAsInt();
 
     // Extract pointer to the underlying ranked descriptor and bitcast it to a
     // memref<element_type> descriptor pointer to minimize the number of GEP
     // operations.
-    UnrankedMemRefDescriptor unrankedDesc(transformed.source());
+    UnrankedMemRefDescriptor unrankedDesc(adaptor.source());
     Value underlyingRankedDesc = unrankedDesc.memRefDescPtr(rewriter, loc);
     Value scalarMemRefDescPtr = rewriter.create<LLVM::BitcastOp>(
         loc,
         LLVM::LLVMPointerType::get(typeConverter->convertType(scalarMemRefType),
                                    addressSpace),
         underlyingRankedDesc);
 
     // Get pointer to offset field of memref<element_type> descriptor.
     Type indexPtrTy = LLVM::LLVMPointerType::get(
         getTypeConverter()->getIndexType(), addressSpace);
     Value two = rewriter.create<LLVM::ConstantOp>(
         loc, typeConverter->convertType(rewriter.getI32Type()),
         rewriter.getI32IntegerAttr(2));
     Value offsetPtr = rewriter.create<LLVM::GEPOp>(
         loc, indexPtrTy, scalarMemRefDescPtr,
         ValueRange({createIndexConstant(rewriter, loc, 0), two}));
 
     // The size value that we have to extract can be obtained using GEPop with
     // `dimOp.index() + 1` index argument.
     Value idxPlusOne = rewriter.create<LLVM::AddOp>(
-        loc, createIndexConstant(rewriter, loc, 1), transformed.index());
+        loc, createIndexConstant(rewriter, loc, 1), adaptor.index());
     Value sizePtr = rewriter.create<LLVM::GEPOp>(loc, indexPtrTy, offsetPtr,
                                                  ValueRange({idxPlusOne}));
     return rewriter.create<LLVM::LoadOp>(loc, sizePtr);
   }
 
   Optional<int64_t> getConstantDimIndex(memref::DimOp dimOp) const {
     if (Optional<int64_t> idx = dimOp.getConstantIndex())
       return idx;
 
     if (auto constantOp = dimOp.index().getDefiningOp<LLVM::ConstantOp>())
       return constantOp.value().cast<IntegerAttr>().getValue().getSExtValue();
 
     return llvm::None;
   }
 
   Value extractSizeOfRankedMemRef(Type operandType, memref::DimOp dimOp,
-                                  ArrayRef<Value> operands,
+                                  OpAdaptor adaptor,
                                   ConversionPatternRewriter &rewriter) const {
     Location loc = dimOp.getLoc();
-    memref::DimOp::Adaptor transformed(operands);
+
     // Take advantage if index is constant.
     MemRefType memRefType = operandType.cast<MemRefType>();
     if (Optional<int64_t> index = getConstantDimIndex(dimOp)) {
       int64_t i = index.getValue();
       if (memRefType.isDynamicDim(i)) {
         // extract dynamic size from the memref descriptor.
-        MemRefDescriptor descriptor(transformed.source());
+        MemRefDescriptor descriptor(adaptor.source());
         return descriptor.size(rewriter, loc, i);
       }
       // Use constant for static size.
       int64_t dimSize = memRefType.getDimSize(i);
       return createIndexConstant(rewriter, loc, dimSize);
     }
-    Value index = transformed.index();
+    Value index = adaptor.index();
     int64_t rank = memRefType.getRank();
-    MemRefDescriptor memrefDescriptor(transformed.source());
+    MemRefDescriptor memrefDescriptor(adaptor.source());
     return memrefDescriptor.size(rewriter, loc, index, rank);
   }
 };
 
 /// Returns the LLVM type of the global variable given the memref type `type`.
 static Type convertGlobalMemrefTypeToLLVM(MemRefType type,
                                           LLVMTypeConverter &typeConverter) {
   // LLVM type for a global memref will be a multi-dimension array. For

 }
 
 /// GlobalMemrefOp is lowered to a LLVM Global Variable.
 struct GlobalMemrefOpLowering
     : public ConvertOpToLLVMPattern<memref::GlobalOp> {
   using ConvertOpToLLVMPattern<memref::GlobalOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(memref::GlobalOp global, ArrayRef<Value> operands,
+  matchAndRewrite(memref::GlobalOp global, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     MemRefType type = global.type();
     if (!isConvertibleAndHasIdentityMaps(type))
       return failure();
 
     Type arrayTy = convertGlobalMemrefTypeToLLVM(type, *getTypeConverter());
 
     LLVM::Linkage linkage =

 };
 
 // Load operation is lowered to obtaining a pointer to the indexed element
 // and loading it.
 struct LoadOpLowering : public LoadStoreOpLowering<memref::LoadOp> {
   using Base::Base;
 
   LogicalResult
-  matchAndRewrite(memref::LoadOp loadOp, ArrayRef<Value> operands,
+  matchAndRewrite(memref::LoadOp loadOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    memref::LoadOp::Adaptor transformed(operands);
     auto type = loadOp.getMemRefType();
 
-    Value dataPtr =
-        getStridedElementPtr(loadOp.getLoc(), type, transformed.memref(),
-                             transformed.indices(), rewriter);
+    Value dataPtr = getStridedElementPtr(
+        loadOp.getLoc(), type, adaptor.memref(), adaptor.indices(), rewriter);
     rewriter.replaceOpWithNewOp<LLVM::LoadOp>(loadOp, dataPtr);
     return success();
   }
 };
 
 // Store operation is lowered to obtaining a pointer to the indexed element,
 // and storing the given value to it.
 struct StoreOpLowering : public LoadStoreOpLowering<memref::StoreOp> {
   using Base::Base;
 
   LogicalResult
-  matchAndRewrite(memref::StoreOp op, ArrayRef<Value> operands,
+  matchAndRewrite(memref::StoreOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     auto type = op.getMemRefType();
-    memref::StoreOp::Adaptor transformed(operands);
 
-    Value dataPtr =
-        getStridedElementPtr(op.getLoc(), type, transformed.memref(),
-                             transformed.indices(), rewriter);
-    rewriter.replaceOpWithNewOp<LLVM::StoreOp>(op, transformed.value(),
-                                               dataPtr);
+    Value dataPtr = getStridedElementPtr(op.getLoc(), type, adaptor.memref(),
+                                         adaptor.indices(), rewriter);
+    rewriter.replaceOpWithNewOp<LLVM::StoreOp>(op, adaptor.value(), dataPtr);
     return success();
   }
 };
 
 // The prefetch operation is lowered in a way similar to the load operation
 // except that the llvm.prefetch operation is used for replacement.
 struct PrefetchOpLowering : public LoadStoreOpLowering<memref::PrefetchOp> {
   using Base::Base;
 
   LogicalResult
-  matchAndRewrite(memref::PrefetchOp prefetchOp, ArrayRef<Value> operands,
+  matchAndRewrite(memref::PrefetchOp prefetchOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    memref::PrefetchOp::Adaptor transformed(operands);
     auto type = prefetchOp.getMemRefType();
     auto loc = prefetchOp.getLoc();
 
-    Value dataPtr = getStridedElementPtr(loc, type, transformed.memref(),
-                                         transformed.indices(), rewriter);
+    Value dataPtr = getStridedElementPtr(loc, type, adaptor.memref(),
+                                         adaptor.indices(), rewriter);
 
     // Replace with llvm.prefetch.
     auto llvmI32Type = typeConverter->convertType(rewriter.getIntegerType(32));
     auto isWrite = rewriter.create<LLVM::ConstantOp>(
         loc, llvmI32Type, rewriter.getI32IntegerAttr(prefetchOp.isWrite()));
     auto localityHint = rewriter.create<LLVM::ConstantOp>(
         loc, llvmI32Type,
         rewriter.getI32IntegerAttr(prefetchOp.localityHint()));
LogicalResult match(memref::CastOp memRefCastOp) const override {
 
     // Unranked to unranked cast is disallowed
     return !(srcType.isa<UnrankedMemRefType>() &&
              dstType.isa<UnrankedMemRefType>())
                ? success()
                : failure();
   }
 
-  void rewrite(memref::CastOp memRefCastOp, ArrayRef<Value> operands,
+  void rewrite(memref::CastOp memRefCastOp, OpAdaptor adaptor,
                ConversionPatternRewriter &rewriter) const override {
-    memref::CastOp::Adaptor transformed(operands);
-
     auto srcType = memRefCastOp.getOperand().getType();
     auto dstType = memRefCastOp.getType();
     auto targetStructType = typeConverter->convertType(memRefCastOp.getType());
     auto loc = memRefCastOp.getLoc();
 
     // For ranked/ranked case, just keep the original descriptor.
     if (srcType.isa<MemRefType>() && dstType.isa<MemRefType>())
-      return rewriter.replaceOp(memRefCastOp, {transformed.source()});
+      return rewriter.replaceOp(memRefCastOp, {adaptor.source()});
 
     if (srcType.isa<MemRefType>() && dstType.isa<UnrankedMemRefType>()) {
       // Casting ranked to unranked memref type
       // Set the rank in the destination from the memref type
       // Allocate space on the stack and copy the src memref descriptor
       // Set the ptr in the destination to the stack space
       auto srcMemRefType = srcType.cast<MemRefType>();
       int64_t rank = srcMemRefType.getRank();
       // ptr = AllocaOp sizeof(MemRefDescriptor)
       auto ptr = getTypeConverter()->promoteOneMemRefDescriptor(
-          loc, transformed.source(), rewriter);
+          loc, adaptor.source(), rewriter);
       // voidptr = BitCastOp srcType* to void*
       auto voidPtr =
           rewriter.create<LLVM::BitcastOp>(loc, getVoidPtrType(), ptr)
               .getResult();
       // rank = ConstantOp srcRank
       auto rankVal = rewriter.create<LLVM::ConstantOp>(
           loc, typeConverter->convertType(rewriter.getIntegerType(64)),
           rewriter.getI64IntegerAttr(rank));
       // undef = UndefOp
       UnrankedMemRefDescriptor memRefDesc =
           UnrankedMemRefDescriptor::undef(rewriter, loc, targetStructType);
       // d1 = InsertValueOp undef, rank, 0
       memRefDesc.setRank(rewriter, loc, rankVal);
       // d2 = InsertValueOp d1, voidptr, 1
       memRefDesc.setMemRefDescPtr(rewriter, loc, voidPtr);
       rewriter.replaceOp(memRefCastOp, (Value)memRefDesc);
 
     } else if (srcType.isa<UnrankedMemRefType>() && dstType.isa<MemRefType>()) {
       // Casting from unranked type to ranked.
       // The operation is assumed to be doing a correct cast. If the destination
       // type mismatches the unranked the type, it is undefined behavior.
-      UnrankedMemRefDescriptor memRefDesc(transformed.source());
+      UnrankedMemRefDescriptor memRefDesc(adaptor.source());
       // ptr = ExtractValueOp src, 1
       auto ptr = memRefDesc.memRefDescPtr(rewriter, loc);
       // castPtr = BitCastOp i8* to structTy*
       auto castPtr =
           rewriter
               .create<LLVM::BitcastOp>(
                   loc, LLVM::LLVMPointerType::get(targetStructType), ptr)
               .getResult();
       // struct = LoadOp castPtr
       auto loadOp = rewriter.create<LLVM::LoadOp>(loc, castPtr);
       rewriter.replaceOp(memRefCastOp, loadOp.getResult());
     } else {
       llvm_unreachable("Unsupported unranked memref to unranked memref cast");
     }
   }
 };
 
 struct MemRefCopyOpLowering : public ConvertOpToLLVMPattern<memref::CopyOp> {
   using ConvertOpToLLVMPattern<memref::CopyOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(memref::CopyOp op, ArrayRef<Value> operands,
+  matchAndRewrite(memref::CopyOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     auto loc = op.getLoc();
-    memref::CopyOp::Adaptor adaptor(operands);
     auto srcType = op.source().getType().cast<BaseMemRefType>();
     auto targetType = op.target().getType().cast<BaseMemRefType>();
 
     // First make sure we have an unranked memref descriptor representation.
     auto makeUnranked = [&, this](Value ranked, BaseMemRefType type) {
       auto rank = rewriter.create<LLVM::ConstantOp>(
           loc, getIndexType(), rewriter.getIndexAttr(type.getRank()));
       auto *typeConverter = getTypeConverter();

 }
 
 struct MemRefReinterpretCastOpLowering
     : public ConvertOpToLLVMPattern<memref::ReinterpretCastOp> {
   using ConvertOpToLLVMPattern<
       memref::ReinterpretCastOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(memref::ReinterpretCastOp castOp, ArrayRef<Value> operands,
+  matchAndRewrite(memref::ReinterpretCastOp castOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    memref::ReinterpretCastOp::Adaptor adaptor(operands,
-                                               castOp->getAttrDictionary());
     Type srcType = castOp.source().getType();
 
     Value descriptor;
     if (failed(convertSourceMemRefToDescriptor(rewriter, srcType, castOp,
                                                adaptor, &descriptor)))
       return failure();
     rewriter.replaceOp(castOp, {descriptor});
     return success();
private:
   }
 };
 
 struct MemRefReshapeOpLowering
     : public ConvertOpToLLVMPattern<memref::ReshapeOp> {
   using ConvertOpToLLVMPattern<memref::ReshapeOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(memref::ReshapeOp reshapeOp, ArrayRef<Value> operands,
+  matchAndRewrite(memref::ReshapeOp reshapeOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    auto *op = reshapeOp.getOperation();
-    memref::ReshapeOp::Adaptor adaptor(operands, op->getAttrDictionary());
     Type srcType = reshapeOp.source().getType();
 
     Value descriptor;
     if (failed(convertSourceMemRefToDescriptor(rewriter, srcType, reshapeOp,
                                                adaptor, &descriptor)))
       return failure();
-    rewriter.replaceOp(op, {descriptor});
+    rewriter.replaceOp(reshapeOp, {descriptor});
     return success();
   }
 
 private:
   LogicalResult
   convertSourceMemRefToDescriptor(ConversionPatternRewriter &rewriter,
                                   Type srcType, memref::ReshapeOp reshapeOp,
                                   memref::ReshapeOp::Adaptor adaptor,

 template <typename ReshapeOp>
 class ReassociatingReshapeOpConversion
     : public ConvertOpToLLVMPattern<ReshapeOp> {
 public:
   using ConvertOpToLLVMPattern<ReshapeOp>::ConvertOpToLLVMPattern;
   using ReshapeOpAdaptor = typename ReshapeOp::Adaptor;
 
   LogicalResult
-  matchAndRewrite(ReshapeOp reshapeOp, ArrayRef<Value> operands,
+  matchAndRewrite(ReshapeOp reshapeOp, typename ReshapeOp::Adaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     MemRefType dstType = reshapeOp.getResultType();
     MemRefType srcType = reshapeOp.getSrcType();
     if (!srcType.getAffineMaps().empty() || !dstType.getAffineMaps().empty()) {
       return rewriter.notifyMatchFailure(reshapeOp,
                                          "only empty layout map is supported");
     }
 
     int64_t offset;
     SmallVector<int64_t, 4> strides;
     if (failed(getStridesAndOffset(dstType, strides, offset))) {
       return rewriter.notifyMatchFailure(
           reshapeOp, "failed to get stride and offset exprs");
     }
 
-    ReshapeOpAdaptor adaptor(operands);
     MemRefDescriptor srcDesc(adaptor.src());
     Location loc = reshapeOp->getLoc();
     auto dstDesc = MemRefDescriptor::undef(
         rewriter, loc, this->typeConverter->convertType(dstType));
     dstDesc.setAllocatedPtr(rewriter, loc, srcDesc.allocatedPtr(rewriter, loc));
     dstDesc.setAlignedPtr(rewriter, loc, srcDesc.alignedPtr(rewriter, loc));
     dstDesc.setOffset(rewriter, loc, srcDesc.offset(rewriter, loc));
 

 ///   1. An `llvm.mlir.undef` operation to create a memref descriptor
 ///   2. Updates to the descriptor to introduce the data ptr, offset, size
 ///      and stride.
 /// The subview op is replaced by the descriptor.
 struct SubViewOpLowering : public ConvertOpToLLVMPattern<memref::SubViewOp> {
   using ConvertOpToLLVMPattern<memref::SubViewOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(memref::SubViewOp subViewOp, ArrayRef<Value> operands,
+  matchAndRewrite(memref::SubViewOp subViewOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     auto loc = subViewOp.getLoc();
 
     auto sourceMemRefType = subViewOp.source().getType().cast<MemRefType>();
     auto sourceElementTy =
         typeConverter->convertType(sourceMemRefType.getElementType());
 
     auto viewMemRefType = subViewOp.getType();
matchAndRewrite(memref::SubViewOp subViewOp, OpAdaptor adaptor,
     // Extract the offset and strides from the type.
     int64_t offset;
     SmallVector<int64_t, 4> strides;
     auto successStrides = getStridesAndOffset(inferredType, strides, offset);
     if (failed(successStrides))
       return failure();
 
     // Create the descriptor.
-    if (!LLVM::isCompatibleType(operands.front().getType()))
+    if (!LLVM::isCompatibleType(adaptor.getOperands().front().getType()))
       return failure();
-    MemRefDescriptor sourceMemRef(operands.front());
+    MemRefDescriptor sourceMemRef(adaptor.getOperands().front());
     auto targetMemRef = MemRefDescriptor::undef(rewriter, loc, targetDescTy);
 
     // Copy the buffer pointer from the old descriptor to the new one.
     Value extracted = sourceMemRef.allocatedPtr(rewriter, loc);
     Value bitcastPtr = rewriter.create<LLVM::BitcastOp>(
         loc,
         LLVM::LLVMPointerType::get(targetElementTy,
                                    viewMemRefType.getMemorySpaceAsInt()),
if (!ShapedType::isDynamicStrideOrOffset(offset)) {
       // because of trailing semantics. In this case, the offset is guaranteed
       // to be interpreted as 0 and we can just skip the extra dimensions.
       for (unsigned i = 0, e = std::min(inferredShapeRank,
                                         subViewOp.getMixedOffsets().size());
            i < e; ++i) {
         Value offset =
             // TODO: need OpFoldResult ODS adaptor to clean this up.
             subViewOp.isDynamicOffset(i)
-                ? operands[subViewOp.getIndexOfDynamicOffset(i)]
+                ? adaptor.getOperands()[subViewOp.getIndexOfDynamicOffset(i)]
                 : rewriter.create<LLVM::ConstantOp>(
                       loc, llvmIndexType,
                       rewriter.getI64IntegerAttr(subViewOp.getStaticOffset(i)));
         Value mul = rewriter.create<LLVM::MulOp>(loc, offset, strideValues[i]);
         baseOffset = rewriter.create<LLVM::AddOp>(loc, baseOffset, mul);
       }
       targetMemRef.setOffset(rewriter, loc, baseOffset);
     }
for (int i = inferredShapeRank - 1, j = resultShapeRank - 1;
           size = cast.getResult(0);
         }
         stride = rewriter.create<LLVM::ConstantOp>(
             loc, llvmIndexType, rewriter.getI64IntegerAttr(1));
       } else {
         // TODO: need OpFoldResult ODS adaptor to clean this up.
         size =
             subViewOp.isDynamicSize(i)
-                ? operands[subViewOp.getIndexOfDynamicSize(i)]
+                ? adaptor.getOperands()[subViewOp.getIndexOfDynamicSize(i)]
                 : rewriter.create<LLVM::ConstantOp>(
                       loc, llvmIndexType,
                       rewriter.getI64IntegerAttr(subViewOp.getStaticSize(i)));
         if (!ShapedType::isDynamicStrideOrOffset(strides[i])) {
           stride = rewriter.create<LLVM::ConstantOp>(
               loc, llvmIndexType, rewriter.getI64IntegerAttr(strides[i]));
         } else {
-          stride = subViewOp.isDynamicStride(i)
-                       ? operands[subViewOp.getIndexOfDynamicStride(i)]
-                       : rewriter.create<LLVM::ConstantOp>(
-                             loc, llvmIndexType,
-                             rewriter.getI64IntegerAttr(
-                                 subViewOp.getStaticStride(i)));
+          stride =
+              subViewOp.isDynamicStride(i)
+                  ? adaptor.getOperands()[subViewOp.getIndexOfDynamicStride(i)]
           stride = rewriter.create<LLVM::MulOp>(loc, stride, strideValues[i]);
         }
       }
       targetMemRef.setSize(rewriter, loc, j, size);
       targetMemRef.setStride(rewriter, loc, j, stride);
       j--;
     }
 

 ///      and stride. Size and stride are permutations of the original values.
 ///   4. A store of the resulting ViewDescriptor to the alloca'ed pointer.
 /// The transpose op is replaced by the alloca'ed pointer.
 class TransposeOpLowering : public ConvertOpToLLVMPattern<memref::TransposeOp> {
 public:
   using ConvertOpToLLVMPattern<memref::TransposeOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(memref::TransposeOp transposeOp, ArrayRef<Value> operands,
+  matchAndRewrite(memref::TransposeOp transposeOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     auto loc = transposeOp.getLoc();
-    memref::TransposeOpAdaptor adaptor(operands);
     MemRefDescriptor viewMemRef(adaptor.in());
 
     // No permutation, early exit.
     if (transposeOp.permutation().isIdentity())
       return rewriter.replaceOp(transposeOp, {viewMemRef}), success();
 
     auto targetMemRef = MemRefDescriptor::undef(
         rewriter, loc, typeConverter->convertType(transposeOp.getShapedType()));
if (nextSize)
       return runningStride
                  ? rewriter.create<LLVM::MulOp>(loc, runningStride, nextSize)
                  : nextSize;
     assert(!runningStride);
     return createIndexConstant(rewriter, loc, 1);
   }
 
   LogicalResult
-  matchAndRewrite(memref::ViewOp viewOp, ArrayRef<Value> operands,
+  matchAndRewrite(memref::ViewOp viewOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     auto loc = viewOp.getLoc();
-    memref::ViewOpAdaptor adaptor(operands);
 
     auto viewMemRefType = viewOp.getType();
     auto targetElementTy =
         typeConverter->convertType(viewMemRefType.getElementType());
     auto targetDescTy = typeConverter->convertType(viewMemRefType);
     if (!targetDescTy || !targetElementTy ||
         !LLVM::isCompatibleType(targetElementTy) ||
         !LLVM::isCompatibleType(targetDescTy))

mlir/lib/Conversion/OpenACCToLLVM/OpenACCToLLVM.cpp

 
 namespace {
 
 template <typename Op>
 class LegalizeDataOpForLLVMTranslation : public ConvertOpToLLVMPattern<Op> {
   using ConvertOpToLLVMPattern<Op>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(Op op, ArrayRef<Value> operands,
+  matchAndRewrite(Op op, typename Op::Adaptor adaptor,
                   ConversionPatternRewriter &builder) const override {
     Location loc = op.getLoc();
     TypeConverter *converter = ConvertToLLVMPattern::getTypeConverter();
 
     unsigned numDataOperand = op.getNumDataOperands();
 
     // Keep the non data operands without modification.
-    auto nonDataOperands =
-        operands.take_front(operands.size() - numDataOperand);
+    auto nonDataOperands = adaptor.getOperands().take_front(
+        adaptor.getOperands().size() - numDataOperand);
     SmallVector<Value> convertedOperands;
     convertedOperands.append(nonDataOperands.begin(), nonDataOperands.end());
 
     // Go over the data operand and legalize them for translation.
     for (unsigned idx = 0; idx < numDataOperand; ++idx) {
       Value originalDataOperand = op.getDataOperand(idx);
 
       // Traverse operands that were converted to MemRefDescriptors.

mlir/lib/Conversion/OpenMPToLLVM/OpenMPToLLVM.cpp

 /// 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 ConvertOpToLLVMPattern<OpType> {
   using ConvertOpToLLVMPattern<OpType>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(OpType curOp, ArrayRef<Value> operands,
+  matchAndRewrite(OpType curOp, typename OpType::Adaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    auto newOp = rewriter.create<OpType>(curOp.getLoc(), TypeRange(), operands,
-                                         curOp->getAttrs());
+    auto newOp = rewriter.create<OpType>(
+        curOp.getLoc(), TypeRange(), adaptor.getOperands(), curOp->getAttrs());
     rewriter.inlineRegionBefore(curOp.region(), newOp.region(),
                                 newOp.region().end());
     if (failed(rewriter.convertRegionTypes(&newOp.region(),
                                            *this->getTypeConverter())))
       return failure();
 
     rewriter.eraseOp(curOp);
     return success();

mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp

 /// This pattern emulates a call to the kernel in LLVM dialect. For that, we
 /// copy the data to the global variable (emulating device side), call the
 /// kernel as a normal void LLVM function, and copy the data back (emulating the
 /// host side).
 class GPULaunchLowering : public ConvertOpToLLVMPattern<gpu::LaunchFuncOp> {
   using ConvertOpToLLVMPattern<gpu::LaunchFuncOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(gpu::LaunchFuncOp launchOp, ArrayRef<Value> operands,
+  matchAndRewrite(gpu::LaunchFuncOp launchOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     auto *op = launchOp.getOperation();
     MLIRContext *context = rewriter.getContext();
     auto module = launchOp->getParentOfType<ModuleOp>();
 
     // Get the SPIR-V module that represents the gpu kernel module. The module
     // is named:
     //   __spv__{kernel_module_name}
matchAndRewrite(gpu::LaunchFuncOp launchOp, OpAdaptor adaptor,
     if (failed(getKernelGlobalVariables(spvModule, globalVariableMap)))
       return failure();
 
     // Traverse kernel operands that were converted to MemRefDescriptors. For
     // each operand, create a global variable and copy data from operand to it.
     Location loc = launchOp.getLoc();
     SmallVector<CopyInfo, 4> copyInfo;
     auto numKernelOperands = launchOp.getNumKernelOperands();
-    auto kernelOperands = operands.take_back(numKernelOperands);
+    auto kernelOperands = adaptor.getOperands().take_back(numKernelOperands);
     for (auto operand : llvm::enumerate(kernelOperands)) {
       // Check if the kernel's operand is a ranked memref.
       auto memRefType = launchOp.getKernelOperand(operand.index())
                             .getType()
                             .dyn_cast<MemRefType>();
       if (!memRefType)
         return failure();
 

mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVM.cpp

 
 /// Conversion pattern for a vector.bitcast.
 class VectorBitCastOpConversion
     : public ConvertOpToLLVMPattern<vector::BitCastOp> {
 public:
   using ConvertOpToLLVMPattern<vector::BitCastOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(vector::BitCastOp bitCastOp, ArrayRef<Value> operands,
+  matchAndRewrite(vector::BitCastOp bitCastOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     // Only 1-D vectors can be lowered to LLVM.
     VectorType resultTy = bitCastOp.getType();
     if (resultTy.getRank() != 1)
       return failure();
     Type newResultTy = typeConverter->convertType(resultTy);
     rewriter.replaceOpWithNewOp<LLVM::BitcastOp>(bitCastOp, newResultTy,
-                                                 operands[0]);
+                                                 adaptor.getOperands()[0]);
     return success();
   }
 };
 
 /// Conversion pattern for a vector.matrix_multiply.
 /// This is lowered directly to the proper llvm.intr.matrix.multiply.
 class VectorMatmulOpConversion
     : public ConvertOpToLLVMPattern<vector::MatmulOp> {
 public:
   using ConvertOpToLLVMPattern<vector::MatmulOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(vector::MatmulOp matmulOp, ArrayRef<Value> operands,
+  matchAndRewrite(vector::MatmulOp matmulOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    auto adaptor = vector::MatmulOpAdaptor(operands);
     rewriter.replaceOpWithNewOp<LLVM::MatrixMultiplyOp>(
         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 ConvertOpToLLVMPattern<vector::FlatTransposeOp> {
 public:
   using ConvertOpToLLVMPattern<vector::FlatTransposeOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(vector::FlatTransposeOp transOp, ArrayRef<Value> operands,
+  matchAndRewrite(vector::FlatTransposeOp transOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    auto adaptor = vector::FlatTransposeOpAdaptor(operands);
     rewriter.replaceOpWithNewOp<LLVM::MatrixTransposeOp>(
         transOp, typeConverter->convertType(transOp.res().getType()),
         adaptor.matrix(), transOp.rows(), transOp.columns());
     return success();
   }
 };
 
 /// Overloaded utility that replaces a vector.load, vector.store,

 /// Conversion pattern for a vector.load, vector.store, vector.maskedload, and
 /// vector.maskedstore.
 template <class LoadOrStoreOp, class LoadOrStoreOpAdaptor>
 class VectorLoadStoreConversion : public ConvertOpToLLVMPattern<LoadOrStoreOp> {
 public:
   using ConvertOpToLLVMPattern<LoadOrStoreOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(LoadOrStoreOp loadOrStoreOp, ArrayRef<Value> operands,
+  matchAndRewrite(LoadOrStoreOp loadOrStoreOp,
+                  typename LoadOrStoreOp::Adaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     // Only 1-D vectors can be lowered to LLVM.
     VectorType vectorTy = loadOrStoreOp.getVectorType();
     if (vectorTy.getRank() > 1)
       return failure();
 
     auto loc = loadOrStoreOp->getLoc();
-    auto adaptor = LoadOrStoreOpAdaptor(operands);
     MemRefType memRefTy = loadOrStoreOp.getMemRefType();
 
     // Resolve alignment.
     unsigned align;
     if (failed(getMemRefOpAlignment(*this->getTypeConverter(), loadOrStoreOp,
                                     align)))
       return failure();
 

 
 /// Conversion pattern for a vector.gather.
 class VectorGatherOpConversion
     : public ConvertOpToLLVMPattern<vector::GatherOp> {
 public:
   using ConvertOpToLLVMPattern<vector::GatherOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(vector::GatherOp gather, ArrayRef<Value> operands,
+  matchAndRewrite(vector::GatherOp gather, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     auto loc = gather->getLoc();
-    auto adaptor = vector::GatherOpAdaptor(operands);
     MemRefType memRefType = gather.getMemRefType();
 
     // Resolve alignment.
     unsigned align;
     if (failed(getMemRefOpAlignment(*getTypeConverter(), gather, align)))
       return failure();
 
     // Resolve address.

 
 /// Conversion pattern for a vector.scatter.
 class VectorScatterOpConversion
     : public ConvertOpToLLVMPattern<vector::ScatterOp> {
 public:
   using ConvertOpToLLVMPattern<vector::ScatterOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(vector::ScatterOp scatter, ArrayRef<Value> operands,
+  matchAndRewrite(vector::ScatterOp scatter, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     auto loc = scatter->getLoc();
-    auto adaptor = vector::ScatterOpAdaptor(operands);
     MemRefType memRefType = scatter.getMemRefType();
 
     // Resolve alignment.
     unsigned align;
     if (failed(getMemRefOpAlignment(*getTypeConverter(), scatter, align)))
       return failure();
 
     // Resolve address.

 
 /// Conversion pattern for a vector.expandload.
 class VectorExpandLoadOpConversion
     : public ConvertOpToLLVMPattern<vector::ExpandLoadOp> {
 public:
   using ConvertOpToLLVMPattern<vector::ExpandLoadOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(vector::ExpandLoadOp expand, ArrayRef<Value> operands,
+  matchAndRewrite(vector::ExpandLoadOp expand, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     auto loc = expand->getLoc();
-    auto adaptor = vector::ExpandLoadOpAdaptor(operands);
     MemRefType memRefType = expand.getMemRefType();
 
     // Resolve address.
     auto vtype = typeConverter->convertType(expand.getVectorType());
     Value ptr = getStridedElementPtr(loc, memRefType, adaptor.base(),
                                      adaptor.indices(), rewriter);
 
     rewriter.replaceOpWithNewOp<LLVM::masked_expandload>(
         expand, vtype, ptr, adaptor.mask(), adaptor.pass_thru());
     return success();
   }
 };
 
 /// Conversion pattern for a vector.compressstore.
 class VectorCompressStoreOpConversion
     : public ConvertOpToLLVMPattern<vector::CompressStoreOp> {
 public:
   using ConvertOpToLLVMPattern<vector::CompressStoreOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(vector::CompressStoreOp compress, ArrayRef<Value> operands,
+  matchAndRewrite(vector::CompressStoreOp compress, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     auto loc = compress->getLoc();
-    auto adaptor = vector::CompressStoreOpAdaptor(operands);
     MemRefType memRefType = compress.getMemRefType();
 
     // Resolve address.
     Value ptr = getStridedElementPtr(loc, memRefType, adaptor.base(),
                                      adaptor.indices(), rewriter);
 
     rewriter.replaceOpWithNewOp<LLVM::masked_compressstore>(
         compress, adaptor.valueToStore(), ptr, adaptor.mask());
     return success();
   }
 };
 
 /// Conversion pattern for all vector reductions.
 class VectorReductionOpConversion
     : public ConvertOpToLLVMPattern<vector::ReductionOp> {
 public:
   explicit VectorReductionOpConversion(LLVMTypeConverter &typeConv,
                                        bool reassociateFPRed)
       : ConvertOpToLLVMPattern<vector::ReductionOp>(typeConv),
         reassociateFPReductions(reassociateFPRed) {}
 
   LogicalResult
-  matchAndRewrite(vector::ReductionOp reductionOp, ArrayRef<Value> operands,
+  matchAndRewrite(vector::ReductionOp reductionOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     auto kind = reductionOp.kind();
     Type eltType = reductionOp.dest().getType();
     Type llvmType = typeConverter->convertType(eltType);
+    Value operand = adaptor.getOperands()[0];
     if (eltType.isIntOrIndex()) {
       // Integer reductions: add/mul/min/max/and/or/xor.
       if (kind == "add")
-        rewriter.replaceOpWithNewOp<LLVM::vector_reduce_add>(
-            reductionOp, llvmType, operands[0]);
+        rewriter.replaceOpWithNewOp<LLVM::vector_reduce_add>(reductionOp,
+                                                             llvmType, operand);
       else if (kind == "mul")
-        rewriter.replaceOpWithNewOp<LLVM::vector_reduce_mul>(
-            reductionOp, llvmType, operands[0]);
+        rewriter.replaceOpWithNewOp<LLVM::vector_reduce_mul>(reductionOp,
+                                                             llvmType, operand);
       else if (kind == "min" &&
                (eltType.isIndex() || eltType.isUnsignedInteger()))
         rewriter.replaceOpWithNewOp<LLVM::vector_reduce_umin>(
-            reductionOp, llvmType, operands[0]);
+            reductionOp, llvmType, operand);
       else if (kind == "min")
         rewriter.replaceOpWithNewOp<LLVM::vector_reduce_smin>(
-            reductionOp, llvmType, operands[0]);
+            reductionOp, llvmType, operand);
       else if (kind == "max" &&
                (eltType.isIndex() || eltType.isUnsignedInteger()))
         rewriter.replaceOpWithNewOp<LLVM::vector_reduce_umax>(
-            reductionOp, llvmType, operands[0]);
+            reductionOp, llvmType, operand);
       else if (kind == "max")
         rewriter.replaceOpWithNewOp<LLVM::vector_reduce_smax>(
-            reductionOp, llvmType, operands[0]);
+            reductionOp, llvmType, operand);
       else if (kind == "and")
-        rewriter.replaceOpWithNewOp<LLVM::vector_reduce_and>(
-            reductionOp, llvmType, operands[0]);
+        rewriter.replaceOpWithNewOp<LLVM::vector_reduce_and>(reductionOp,
+                                                             llvmType, operand);
       else if (kind == "or")
-        rewriter.replaceOpWithNewOp<LLVM::vector_reduce_or>(
-            reductionOp, llvmType, operands[0]);
+        rewriter.replaceOpWithNewOp<LLVM::vector_reduce_or>(reductionOp,
+                                                            llvmType, operand);
       else if (kind == "xor")
-        rewriter.replaceOpWithNewOp<LLVM::vector_reduce_xor>(
-            reductionOp, llvmType, operands[0]);
+        rewriter.replaceOpWithNewOp<LLVM::vector_reduce_xor>(reductionOp,
+                                                             llvmType, operand);
       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>(
-                                            reductionOp->getLoc(), llvmType,
-                                            rewriter.getZeroAttr(eltType));
+      Value acc = adaptor.getOperands().size() > 1
+                      ? adaptor.getOperands()[1]
       rewriter.replaceOpWithNewOp<LLVM::vector_reduce_fadd>(
-          reductionOp, llvmType, acc, operands[0],
+          reductionOp, llvmType, acc, operand,
           rewriter.getBoolAttr(reassociateFPReductions));
     } else if (kind == "mul") {
       // Optional accumulator (or one).
-      Value acc = operands.size() > 1
-                      ? operands[1]
+      Value acc = adaptor.getOperands().size() > 1
+                      ? adaptor.getOperands()[1]
                       : rewriter.create<LLVM::ConstantOp>(
                             reductionOp->getLoc(), llvmType,
                             rewriter.getFloatAttr(eltType, 1.0));
       rewriter.replaceOpWithNewOp<LLVM::vector_reduce_fmul>(
-          reductionOp, llvmType, acc, operands[0],
+          reductionOp, llvmType, acc, operand,
           rewriter.getBoolAttr(reassociateFPReductions));
     } else if (kind == "min")
-      rewriter.replaceOpWithNewOp<LLVM::vector_reduce_fmin>(
-          reductionOp, llvmType, operands[0]);
+      rewriter.replaceOpWithNewOp<LLVM::vector_reduce_fmin>(reductionOp,
+                                                            llvmType, operand);
     else if (kind == "max")
-      rewriter.replaceOpWithNewOp<LLVM::vector_reduce_fmax>(
-          reductionOp, llvmType, operands[0]);
+      rewriter.replaceOpWithNewOp<LLVM::vector_reduce_fmax>(reductionOp,
+                                                            llvmType, operand);
     else
       return failure();
     return success();
   }
 
 private:
   const bool reassociateFPReductions;
 };
 
 class VectorShuffleOpConversion
     : public ConvertOpToLLVMPattern<vector::ShuffleOp> {
 public:
   using ConvertOpToLLVMPattern<vector::ShuffleOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(vector::ShuffleOp shuffleOp, ArrayRef<Value> operands,
+  matchAndRewrite(vector::ShuffleOp shuffleOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     auto loc = shuffleOp->getLoc();
-    auto adaptor = vector::ShuffleOpAdaptor(operands);
     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)

 
 class VectorExtractElementOpConversion
     : public ConvertOpToLLVMPattern<vector::ExtractElementOp> {
 public:
   using ConvertOpToLLVMPattern<
       vector::ExtractElementOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(vector::ExtractElementOp extractEltOp,
-                  ArrayRef<Value> operands,
+  matchAndRewrite(vector::ExtractElementOp extractEltOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    auto adaptor = vector::ExtractElementOpAdaptor(operands);
     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>(
         extractEltOp, llvmType, adaptor.vector(), adaptor.position());
     return success();
   }
 };
 
 class VectorExtractOpConversion
     : public ConvertOpToLLVMPattern<vector::ExtractOp> {
 public:
   using ConvertOpToLLVMPattern<vector::ExtractOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(vector::ExtractOp extractOp, ArrayRef<Value> operands,
+  matchAndRewrite(vector::ExtractOp extractOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     auto loc = extractOp->getLoc();
-    auto adaptor = vector::ExtractOpAdaptor(operands);
     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();

 ///    (!llvm."<8 x f32>">, !llvm<"<8 x f32>">, !llvm<"<8 x f32>">)
 ///    -> !llvm."<8 x f32>">
 /// ```
 class VectorFMAOp1DConversion : public ConvertOpToLLVMPattern<vector::FMAOp> {
 public:
   using ConvertOpToLLVMPattern<vector::FMAOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(vector::FMAOp fmaOp, ArrayRef<Value> operands,
+  matchAndRewrite(vector::FMAOp fmaOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    auto adaptor = vector::FMAOpAdaptor(operands);
     VectorType vType = fmaOp.getVectorType();
     if (vType.getRank() != 1)
       return failure();
     rewriter.replaceOpWithNewOp<LLVM::FMulAddOp>(fmaOp, adaptor.lhs(),
                                                  adaptor.rhs(), adaptor.acc());
     return success();
   }
 };
 
 class VectorInsertElementOpConversion
     : public ConvertOpToLLVMPattern<vector::InsertElementOp> {
 public:
   using ConvertOpToLLVMPattern<vector::InsertElementOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(vector::InsertElementOp insertEltOp, ArrayRef<Value> operands,
+  matchAndRewrite(vector::InsertElementOp insertEltOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    auto adaptor = vector::InsertElementOpAdaptor(operands);
     auto vectorType = insertEltOp.getDestVectorType();
     auto llvmType = typeConverter->convertType(vectorType);
 
     // Bail if result type cannot be lowered.
     if (!llvmType)
       return failure();
 
     rewriter.replaceOpWithNewOp<LLVM::InsertElementOp>(
         insertEltOp, llvmType, adaptor.dest(), adaptor.source(),
         adaptor.position());
     return success();
   }
 };
 
 class VectorInsertOpConversion
     : public ConvertOpToLLVMPattern<vector::InsertOp> {
 public:
   using ConvertOpToLLVMPattern<vector::InsertOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(vector::InsertOp insertOp, ArrayRef<Value> operands,
+  matchAndRewrite(vector::InsertOp insertOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     auto loc = insertOp->getLoc();
-    auto adaptor = vector::InsertOpAdaptor(operands);
     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();

 }
 
 class VectorTypeCastOpConversion
     : public ConvertOpToLLVMPattern<vector::TypeCastOp> {
 public:
   using ConvertOpToLLVMPattern<vector::TypeCastOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(vector::TypeCastOp castOp, ArrayRef<Value> operands,
+  matchAndRewrite(vector::TypeCastOp castOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     auto loc = castOp->getLoc();
     MemRefType sourceMemRefType =
         castOp.getOperand().getType().cast<MemRefType>();
     MemRefType targetMemRefType = castOp.getType();
 
     // Only static shape casts supported atm.
     if (!sourceMemRefType.hasStaticShape() ||
         !targetMemRefType.hasStaticShape())
       return failure();
 
     auto llvmSourceDescriptorTy =
-        operands[0].getType().dyn_cast<LLVM::LLVMStructType>();
+        adaptor.getOperands()[0].getType().dyn_cast<LLVM::LLVMStructType>();
     if (!llvmSourceDescriptorTy)
       return failure();
-    MemRefDescriptor sourceMemRef(operands[0]);
+    MemRefDescriptor sourceMemRef(adaptor.getOperands()[0]);
 
     auto llvmTargetDescriptorTy = typeConverter->convertType(targetMemRefType)
                                       .dyn_cast_or_null<LLVM::LLVMStructType>();
     if (!llvmTargetDescriptorTy)
       return failure();
 
     // Only contiguous source buffers supported atm.
     auto sourceStrides = computeContiguousStrides(sourceMemRefType);
public:
   // 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(vector::PrintOp printOp, ArrayRef<Value> operands,
+  matchAndRewrite(vector::PrintOp printOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    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>();

mlir/lib/Conversion/VectorToROCDL/VectorToROCDL.cpp

 #include "mlir/Dialect/Vector/VectorOps.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Transforms/DialectConversion.h"
 
 using namespace mlir;
 using namespace mlir::vector;
 
 static LogicalResult replaceTransferOpWithMubuf(
-    ConversionPatternRewriter &rewriter, ArrayRef<Value> operands,
+    ConversionPatternRewriter &rewriter, ValueRange operands,
     LLVMTypeConverter &typeConverter, Location loc, TransferReadOp xferOp,
     Type &vecTy, Value &dwordConfig, Value &vindex, Value &offsetSizeInBytes,
     Value &glc, Value &slc) {
   rewriter.replaceOpWithNewOp<ROCDL::MubufLoadOp>(
       xferOp, vecTy, dwordConfig, vindex, offsetSizeInBytes, glc, slc);
   return success();
 }
 
 static LogicalResult replaceTransferOpWithMubuf(
-    ConversionPatternRewriter &rewriter, ArrayRef<Value> operands,
+    ConversionPatternRewriter &rewriter, ValueRange operands,
     LLVMTypeConverter &typeConverter, Location loc, TransferWriteOp xferOp,
     Type &vecTy, Value &dwordConfig, Value &vindex, Value &offsetSizeInBytes,
     Value &glc, Value &slc) {
   auto adaptor = TransferWriteOpAdaptor(operands, xferOp->getAttrDictionary());
   rewriter.replaceOpWithNewOp<ROCDL::MubufStoreOp>(xferOp, adaptor.vector(),
                                                    dwordConfig, vindex,
                                                    offsetSizeInBytes, glc, slc);
   return success();
 }
 
 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 ConvertOpToLLVMPattern<ConcreteOp> {
 public:
   using ConvertOpToLLVMPattern<ConcreteOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(ConcreteOp xferOp, ArrayRef<Value> operands,
+  matchAndRewrite(ConcreteOp xferOp, typename ConcreteOp::Adaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    typename ConcreteOp::Adaptor adaptor(operands, xferOp->getAttrDictionary());
-
     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.
matchAndRewrite(ConcreteOp xferOp, typename ConcreteOp::Adaptor adaptor,
     // 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, *this->getTypeConverter(), loc, xferOp, vecTy,
-        dwordConfig, int32Zero, int32Zero, int1False, int1False);
+        rewriter, adaptor.getOperands(), *this->getTypeConverter(), loc, xferOp,
+        vecTy, dwordConfig, int32Zero, int32Zero, int1False, int1False);
   }
 };
 } // end anonymous namespace
 
 void mlir::populateVectorToROCDLConversionPatterns(
     LLVMTypeConverter &converter, RewritePatternSet &patterns) {
   patterns.add<VectorTransferConversion<TransferReadOp>,
                VectorTransferConversion<TransferWriteOp>>(converter);

mlir/lib/Dialect/Linalg/Transforms/Bufferize.cpp

 /// This pattern is arguable a std pattern once linalg::CopyOp becomes
 /// std::CopyOp.
 class InsertSliceOpConverter
     : public OpConversionPattern<tensor::InsertSliceOp> {
 public:
   using OpConversionPattern<tensor::InsertSliceOp>::OpConversionPattern;
 
   LogicalResult
-  matchAndRewrite(tensor::InsertSliceOp op, ArrayRef<Value> operands,
+  matchAndRewrite(tensor::InsertSliceOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const final {
-    tensor::InsertSliceOpAdaptor adaptor(operands, op->getAttrDictionary());
     Value sourceMemRef = adaptor.source();
     assert(sourceMemRef.getType().isa<MemRefType>());
 
     // For now, be conservative and copy the converted input memref.
     // In general, the converted input memref here could be aliased or could
     // point into constant memory, so mutating it would lead to miscompilations.
     Value destMemRef = cloneMemref(op.getLoc(), adaptor.dest(), rewriter);
     assert(destMemRef.getType().isa<MemRefType>());

 };
 
 class VectorTransferReadOpConverter
     : public OpConversionPattern<vector::TransferReadOp> {
 public:
   using OpConversionPattern<vector::TransferReadOp>::OpConversionPattern;
 
   LogicalResult
-  matchAndRewrite(vector::TransferReadOp readOp, ArrayRef<Value> operands,
+  matchAndRewrite(vector::TransferReadOp readOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const final {
     if (readOp.getShapedType().isa<MemRefType>())
       return failure();
-    vector::TransferReadOp::Adaptor adaptor(operands,
-                                            readOp->getAttrDictionary());
     rewriter.replaceOpWithNewOp<vector::TransferReadOp>(
         readOp, readOp.getType(), adaptor.source(), adaptor.indices(),
         adaptor.permutation_map(), adaptor.padding(), adaptor.mask(),
         adaptor.in_bounds());
     return success();
   }
 };
 
 class VectorTransferWriteOpConverter
     : public OpConversionPattern<vector::TransferWriteOp> {
 public:
   using OpConversionPattern<vector::TransferWriteOp>::OpConversionPattern;
 
   LogicalResult
-  matchAndRewrite(vector::TransferWriteOp writeOp, ArrayRef<Value> operands,
+  matchAndRewrite(vector::TransferWriteOp writeOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const final {
     if (writeOp.getShapedType().isa<MemRefType>())
       return failure();
-    vector::TransferWriteOp::Adaptor adaptor(operands,
-                                             writeOp->getAttrDictionary());
     rewriter.create<vector::TransferWriteOp>(
         writeOp.getLoc(), adaptor.vector(), adaptor.source(), adaptor.indices(),
         adaptor.permutation_map(),
         adaptor.in_bounds() ? adaptor.in_bounds() : ArrayAttr());
     rewriter.replaceOp(writeOp, adaptor.source());
     return success();
   }
 };

mlir/test/lib/Conversion/StandardToLLVM/TestConvertCallOp.cpp

 
 class TestTypeProducerOpConverter
     : public ConvertOpToLLVMPattern<test::TestTypeProducerOp> {
 public:
   using ConvertOpToLLVMPattern<
       test::TestTypeProducerOp>::ConvertOpToLLVMPattern;
 
   LogicalResult
-  matchAndRewrite(test::TestTypeProducerOp op, ArrayRef<Value> operands,
+  matchAndRewrite(test::TestTypeProducerOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     rewriter.replaceOpWithNewOp<LLVM::NullOp>(op, getVoidPtrType());
     return success();
   }
 };
 
 class TestConvertCallOp
     : public PassWrapper<TestConvertCallOp, OperationPass<ModuleOp>> {

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

 ///   %1 = test.one_variadic_out_one_variadic_in1"(%0)
 /// is replaced with:
 ///   %1 = test.one_variadic_out_one_variadic_in1"(%0, %0)
 struct OneVResOneVOperandOp1Converter
     : public OpConversionPattern<OneVResOneVOperandOp1> {
   using OpConversionPattern<OneVResOneVOperandOp1>::OpConversionPattern;
 
   LogicalResult
-  matchAndRewrite(OneVResOneVOperandOp1 op, ArrayRef<Value> operands,
+  matchAndRewrite(OneVResOneVOperandOp1 op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     auto origOps = op.getOperands();
     assert(std::distance(origOps.begin(), origOps.end()) == 1 &&
            "One operand expected");
     Value origOp = *origOps.begin();
     SmallVector<Value, 2> remappedOperands;
     // Replicate the remapped original operand twice. Note that we don't used
     // the remapped 'operand' since the goal is testing 'getRemappedValue'.

 // Test type conversions
 //===----------------------------------------------------------------------===//
 
 namespace {
 struct TestTypeConversionProducer
     : public OpConversionPattern<TestTypeProducerOp> {
   using OpConversionPattern<TestTypeProducerOp>::OpConversionPattern;
   LogicalResult
-  matchAndRewrite(TestTypeProducerOp op, ArrayRef<Value> operands,
+  matchAndRewrite(TestTypeProducerOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const final {
     Type resultType = op.getType();
     if (resultType.isa<FloatType>())
       resultType = rewriter.getF64Type();
     else if (resultType.isInteger(16))
       resultType = rewriter.getIntegerType(64);
     else
       return failure();
 
     rewriter.replaceOpWithNewOp<TestTypeProducerOp>(op, resultType);
     return success();
   }
 };
 
 /// Call signature conversion and then fail the rewrite to trigger the undo
 /// mechanism.
 struct TestSignatureConversionUndo
     : public OpConversionPattern<TestSignatureConversionUndoOp> {
   using OpConversionPattern<TestSignatureConversionUndoOp>::OpConversionPattern;
 
   LogicalResult
-  matchAndRewrite(TestSignatureConversionUndoOp op, ArrayRef<Value> operands,
+  matchAndRewrite(TestSignatureConversionUndoOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const final {
     (void)rewriter.convertRegionTypes(&op->getRegion(0), *getTypeConverter());
     return failure();
   }
 };
 
 /// 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;
 
   LogicalResult
-  matchAndRewrite(TestTypeConsumerOp op, ArrayRef<Value> operands,
+  matchAndRewrite(TestTypeConsumerOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const final {
-    rewriter.updateRootInPlace(op, [&] { op->setOperands(operands); });
+    rewriter.updateRootInPlace(op,
+                               [&] { op->setOperands(adaptor.getOperands()); });
     return success();
   }
 };
 
 struct TestTypeConversionAnotherProducer
     : public OpRewritePattern<TestAnotherTypeProducerOp> {
   using OpRewritePattern<TestAnotherTypeProducerOp>::OpRewritePattern;
 

 //===----------------------------------------------------------------------===//
 
 namespace {
 /// A rewriter pattern that tests that blocks can be merged.
 struct TestMergeBlock : public OpConversionPattern<TestMergeBlocksOp> {
   using OpConversionPattern<TestMergeBlocksOp>::OpConversionPattern;
 
   LogicalResult
-  matchAndRewrite(TestMergeBlocksOp op, ArrayRef<Value> operands,
+  matchAndRewrite(TestMergeBlocksOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const final {
     Block &firstBlock = op.body().front();
     Operation *branchOp = firstBlock.getTerminator();
     Block *secondBlock = &*(std::next(op.body().begin()));
     auto succOperands = branchOp->getOperands();
     SmallVector<Value, 2> replacements(succOperands);
     rewriter.eraseOp(branchOp);
     rewriter.mergeBlocks(secondBlock, &firstBlock, replacements);

 /// A rewrite mechanism to inline the body of the op into its parent, when both
 /// ops can have a single block.
 struct TestMergeSingleBlockOps
     : public OpConversionPattern<SingleBlockImplicitTerminatorOp> {
   using OpConversionPattern<
       SingleBlockImplicitTerminatorOp>::OpConversionPattern;
 
   LogicalResult
-  matchAndRewrite(SingleBlockImplicitTerminatorOp op, ArrayRef<Value> operands,
+  matchAndRewrite(SingleBlockImplicitTerminatorOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const final {
     SingleBlockImplicitTerminatorOp parentOp =
         op->getParentOfType<SingleBlockImplicitTerminatorOp>();
     if (!parentOp)
       return failure();
     Block &innerBlock = op.region().front();
     TerminatorOp innerTerminator =
         cast<TerminatorOp>(innerBlock.getTerminator());