[mlir][NFC] Replace all usages of PatternMatchResult with LogicalResult

[rriddle]

Summary

This also replaces usages of matchSuccess/matchFailure with success/failure respectively.

Depends On D76203

mlir/docs/DialectConversion.md

 a special `PatternRewriter`, `ConversionPatternRewriter`, that provides special
 hooks for use with the conversion infrastructure.
 
 ```c++
 struct MyConversionPattern : public ConversionPattern {
   /// The `matchAndRewrite` hooks on ConversionPatterns take an additional
   /// `operands` parameter, containing the remapped operands of the original
   /// operation.
-  virtual PatternMatchResult
+  virtual LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const;
 };
 ```
 
 These patterns have the same [restrictions](#restrictions) as the basic rewrite
 patterns used in dialect conversion.
 

mlir/docs/QuickstartRewrites.md

 
 ```c++
 /// Multi-step rewrite using "match" and "rewrite". This allows for separating
 /// the concerns of matching and rewriting.
 struct ConvertTFLeakyRelu : public RewritePattern {
   ConvertTFLeakyRelu(MLIRContext *context)
       : RewritePattern("tf.LeakyRelu", 1, context) {}
 
-  PatternMatchResult match(Operation *op) const override {
-    return matchSuccess();
+  LogicalResult match(Operation *op) const override {
+    return success();
   }
 
   void rewrite(Operation *op, PatternRewriter &rewriter) const override {
     rewriter.replaceOpWithNewOp<TFL::LeakyReluOp>(
         op, op->getResult(0).getType(), op->getOperand(0),
         /*alpha=*/op->getAttrOfType<FloatAttr>("alpha"));
   }
 };
 
 /// Single-step rewrite with "matchAndRewrite". This allows for performing the
 /// rewrite immediately upon a successful match.
 struct ConvertTFLeakyRelu : public RewritePattern {
   ConvertTFLeakyRelu(MLIRContext *context)
       : RewritePattern("tf.LeakyRelu", 1, context) {}
 
-  PatternMatchResult matchAndRewrite(Operation *op,
+  LogicalResult matchAndRewrite(Operation *op,
                                      PatternRewriter &rewriter) const override {
     rewriter.replaceOpWithNewOp<TFL::LeakyReluOp>(
         op, op->getResult(0).getType(), op->getOperand(0),
         /*alpha=*/op->getAttrOfType<FloatAttr>("alpha"));
-    return matchSuccess();
+    return success();
   }
 };
 ```
 
 In the C++ rewrite the static benefit of the rewrite pattern is specified at
 construction. While in the pattern generator a simple heuristic is currently
 employed based around the number of ops matched and replaced.
 

mlir/docs/Tutorials/Toy/Ch-3.md

struct SimplifyRedundantTranspose : public mlir::OpRewritePattern<TransposeOp> {
   /// The "benefit" is used by the framework to order the patterns and process
   /// them in order of profitability.
   SimplifyRedundantTranspose(mlir::MLIRContext *context)
       : OpRewritePattern<TransposeOp>(context, /*benefit=*/1) {}
 
   /// This method is attempting to match a pattern and rewrite it. The rewriter
   /// argument is the orchestrator of the sequence of rewrites. It is expected
   /// to interact with it to perform any changes to the IR from here.
-  mlir::PatternMatchResult
+  mlir::LogicalResult
   matchAndRewrite(TransposeOp op,
                   mlir::PatternRewriter &rewriter) const override {
     // Look through the input of the current transpose.
     mlir::Value transposeInput = op.getOperand();
     TransposeOp transposeInputOp =
         llvm::dyn_cast_or_null<TransposeOp>(transposeInput.getDefiningOp());
 
     // Input defined by another transpose? If not, no match.
     if (!transposeInputOp)
-      return matchFailure();
+      return failure();
 
     // Otherwise, we have a redundant transpose. Use the rewriter.
     rewriter.replaceOp(op, {transposeInputOp.getOperand()}, {transposeInputOp});
-    return matchSuccess();
+    return success();
   }
 };
 ```
 
 The implementation of this rewriter is in `ToyCombine.cpp`. The
 [canonicalization pass](../../Canonicalization.md) applies transformations
 defined by operations in a greedy, iterative manner. To ensure that the
 canonicalization pass applies our new transform, we set

mlir/docs/Tutorials/Toy/Ch-5.md

 ```c++
 /// Lower the `toy.transpose` operation to an affine loop nest.
 struct TransposeOpLowering : public mlir::ConversionPattern {
   TransposeOpLowering(mlir::MLIRContext *ctx)
       : mlir::ConversionPattern(TransposeOp::getOperationName(), 1, ctx) {}
 
   /// Match and rewrite the given `toy.transpose` operation, with the given
   /// operands that have been remapped from `tensor<...>` to `memref<...>`.
-  mlir::PatternMatchResult
+  mlir::LogicalResult
   matchAndRewrite(mlir::Operation *op, ArrayRef<mlir::Value> operands,
                   mlir::ConversionPatternRewriter &rewriter) const final {
     auto loc = op->getLoc();
 
     // Call to a helper function that will lower the current operation to a set
     // of affine loops. We provide a functor that operates on the remapped
     // operands, as well as the loop induction variables for the inner most
     // loop body.
lowerOpToLoops(
           TransposeOpOperandAdaptor transposeAdaptor(memRefOperands);
           mlir::Value input = transposeAdaptor.input();
 
           // Transpose the elements by generating a load from the reverse
           // indices.
           SmallVector<mlir::Value, 2> reverseIvs(llvm::reverse(loopIvs));
           return rewriter.create<mlir::AffineLoadOp>(loc, input, reverseIvs);
         });
-    return matchSuccess();
+    return success();
   }
 };
 ```
 
 Now we can prepare the list of patterns to use during the lowering process:
 
 ```c++
 void ToyToAffineLoweringPass::runOnFunction() {

mlir/examples/toy/Ch3/mlir/ToyCombine.cpp

struct SimplifyRedundantTranspose : public mlir::OpRewritePattern<TransposeOp> {
   /// The "benefit" is used by the framework to order the patterns and process
   /// them in order of profitability.
   SimplifyRedundantTranspose(mlir::MLIRContext *context)
       : OpRewritePattern<TransposeOp>(context, /*benefit=*/1) {}
 
   /// This method attempts to match a pattern and rewrite it. The rewriter
   /// argument is the orchestrator of the sequence of rewrites. The pattern is
   /// expected to interact with it to perform any changes to the IR from here.
-  mlir::PatternMatchResult
+  mlir::LogicalResult
   matchAndRewrite(TransposeOp op,
                   mlir::PatternRewriter &rewriter) const override {
     // Look through the input of the current transpose.
     mlir::Value transposeInput = op.getOperand();
     TransposeOp transposeInputOp =
         llvm::dyn_cast_or_null<TransposeOp>(transposeInput.getDefiningOp());
 
     // Input defined by another transpose? If not, no match.
     if (!transposeInputOp)
-      return matchFailure();
+      return failure();
 
     // Otherwise, we have a redundant transpose. Use the rewriter.
     rewriter.replaceOp(op, {transposeInputOp.getOperand()});
-    return matchSuccess();
+    return success();
   }
 };
 
 /// Register our patterns as "canonicalization" patterns on the TransposeOp so
 /// that they can be picked up by the Canonicalization framework.
 void TransposeOp::getCanonicalizationPatterns(OwningRewritePatternList &results,
                                               MLIRContext *context) {
   results.insert<SimplifyRedundantTranspose>(context);

mlir/examples/toy/Ch4/mlir/ToyCombine.cpp

struct SimplifyRedundantTranspose : public mlir::OpRewritePattern<TransposeOp> {
   /// The "benefit" is used by the framework to order the patterns and process
   /// them in order of profitability.
   SimplifyRedundantTranspose(mlir::MLIRContext *context)
       : OpRewritePattern<TransposeOp>(context, /*benefit=*/1) {}
 
   /// This method attempts to match a pattern and rewrite it. The rewriter
   /// argument is the orchestrator of the sequence of rewrites. The pattern is
   /// expected to interact with it to perform any changes to the IR from here.
-  mlir::PatternMatchResult
+  mlir::LogicalResult
   matchAndRewrite(TransposeOp op,
                   mlir::PatternRewriter &rewriter) const override {
     // Look through the input of the current transpose.
     mlir::Value transposeInput = op.getOperand();
     TransposeOp transposeInputOp =
         llvm::dyn_cast_or_null<TransposeOp>(transposeInput.getDefiningOp());
 
     // Input defined by another transpose? If not, no match.
     if (!transposeInputOp)
-      return matchFailure();
+      return failure();
 
     // Otherwise, we have a redundant transpose. Use the rewriter.
     rewriter.replaceOp(op, {transposeInputOp.getOperand()});
-    return matchSuccess();
+    return success();
   }
 };
 
 /// Register our patterns as "canonicalization" patterns on the TransposeOp so
 /// that they can be picked up by the Canonicalization framework.
 void TransposeOp::getCanonicalizationPatterns(OwningRewritePatternList &results,
                                               MLIRContext *context) {
   results.insert<SimplifyRedundantTranspose>(context);

mlir/examples/toy/Ch5/mlir/LowerToAffineLoops.cpp

 // ToyToAffine RewritePatterns: Binary operations
 //===----------------------------------------------------------------------===//
 
 template <typename BinaryOp, typename LoweredBinaryOp>
 struct BinaryOpLowering : public ConversionPattern {
   BinaryOpLowering(MLIRContext *ctx)
       : ConversionPattern(BinaryOp::getOperationName(), 1, ctx) {}
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const final {
     auto loc = op->getLoc();
     lowerOpToLoops(
         op, operands, rewriter,
         [loc](PatternRewriter &rewriter, ArrayRef<Value> memRefOperands,
               ArrayRef<Value> loopIvs) {
           // Generate an adaptor for the remapped operands of the BinaryOp. This
           // allows for using the nice named accessors that are generated by the
           // ODS.
           typename BinaryOp::OperandAdaptor binaryAdaptor(memRefOperands);
 
           // Generate loads for the element of 'lhs' and 'rhs' at the inner
           // loop.
           auto loadedLhs =
               rewriter.create<AffineLoadOp>(loc, binaryAdaptor.lhs(), loopIvs);
           auto loadedRhs =
               rewriter.create<AffineLoadOp>(loc, binaryAdaptor.rhs(), loopIvs);
 
           // Create the binary operation performed on the loaded values.
           return rewriter.create<LoweredBinaryOp>(loc, loadedLhs, loadedRhs);
         });
-    return matchSuccess();
+    return success();
   }
 };
 using AddOpLowering = BinaryOpLowering<toy::AddOp, AddFOp>;
 using MulOpLowering = BinaryOpLowering<toy::MulOp, MulFOp>;
 
 //===----------------------------------------------------------------------===//
 // ToyToAffine RewritePatterns: Constant operations
 //===----------------------------------------------------------------------===//
 
 struct ConstantOpLowering : public OpRewritePattern<toy::ConstantOp> {
   using OpRewritePattern<toy::ConstantOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(toy::ConstantOp op,
-                                     PatternRewriter &rewriter) const final {
+  LogicalResult matchAndRewrite(toy::ConstantOp op,
     DenseElementsAttr constantValue = op.value();
     Location loc = op.getLoc();
 
     // When lowering the constant operation, we allocate and assign the constant
     // values to a corresponding memref allocation.
     auto tensorType = op.getType().cast<TensorType>();
     auto memRefType = convertTensorToMemRef(tensorType);
     auto alloc = insertAllocAndDealloc(memRefType, loc, rewriter);
std::function<void(uint64_t)> storeElements = [&](uint64_t dimension) {
       }
     };
 
     // Start the element storing recursion from the first dimension.
     storeElements(/*dimension=*/0);
 
     // Replace this operation with the generated alloc.
     rewriter.replaceOp(op, alloc);
-    return matchSuccess();
+    return success();
   }
 };
 
 //===----------------------------------------------------------------------===//
 // ToyToAffine RewritePatterns: Return operations
 //===----------------------------------------------------------------------===//
 
 struct ReturnOpLowering : public OpRewritePattern<toy::ReturnOp> {
   using OpRewritePattern<toy::ReturnOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(toy::ReturnOp op,
-                                     PatternRewriter &rewriter) const final {
+  LogicalResult matchAndRewrite(toy::ReturnOp op,
     // During this lowering, we expect that all function calls have been
     // inlined.
     if (op.hasOperand())
-      return matchFailure();
+      return failure();
 
     // We lower "toy.return" directly to "std.return".
     rewriter.replaceOpWithNewOp<ReturnOp>(op);
-    return matchSuccess();
+    return success();
   }
 };
 
 //===----------------------------------------------------------------------===//
 // ToyToAffine RewritePatterns: Transpose operations
 //===----------------------------------------------------------------------===//
 
 struct TransposeOpLowering : public ConversionPattern {
   TransposeOpLowering(MLIRContext *ctx)
       : ConversionPattern(toy::TransposeOp::getOperationName(), 1, ctx) {}
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const final {
     auto loc = op->getLoc();
     lowerOpToLoops(
         op, operands, rewriter,
         [loc](PatternRewriter &rewriter, ArrayRef<Value> memRefOperands,
               ArrayRef<Value> loopIvs) {
           // Generate an adaptor for the remapped operands of the TransposeOp.
           // This allows for using the nice named accessors that are generated
           // by the ODS.
           toy::TransposeOpOperandAdaptor transposeAdaptor(memRefOperands);
           Value input = transposeAdaptor.input();
 
           // Transpose the elements by generating a load from the reverse
           // indices.
           SmallVector<Value, 2> reverseIvs(llvm::reverse(loopIvs));
           return rewriter.create<AffineLoadOp>(loc, input, reverseIvs);
         });
-    return matchSuccess();
+    return success();
   }
 };
 
 } // end anonymous namespace.
 
 //===----------------------------------------------------------------------===//
 // ToyToAffineLoweringPass
 //===----------------------------------------------------------------------===//

mlir/examples/toy/Ch5/mlir/ToyCombine.cpp

struct SimplifyRedundantTranspose : public mlir::OpRewritePattern<TransposeOp> {
   /// The "benefit" is used by the framework to order the patterns and process
   /// them in order of profitability.
   SimplifyRedundantTranspose(mlir::MLIRContext *context)
       : OpRewritePattern<TransposeOp>(context, /*benefit=*/1) {}
 
   /// This method attempts to match a pattern and rewrite it. The rewriter
   /// argument is the orchestrator of the sequence of rewrites. The pattern is
   /// expected to interact with it to perform any changes to the IR from here.
-  mlir::PatternMatchResult
+  mlir::LogicalResult
   matchAndRewrite(TransposeOp op,
                   mlir::PatternRewriter &rewriter) const override {
     // Look through the input of the current transpose.
     mlir::Value transposeInput = op.getOperand();
     TransposeOp transposeInputOp =
         llvm::dyn_cast_or_null<TransposeOp>(transposeInput.getDefiningOp());
 
     // Input defined by another transpose? If not, no match.
     if (!transposeInputOp)
-      return matchFailure();
+      return failure();
 
     // Otherwise, we have a redundant transpose. Use the rewriter.
     rewriter.replaceOp(op, {transposeInputOp.getOperand()});
-    return matchSuccess();
+    return success();
   }
 };
 
 /// Register our patterns as "canonicalization" patterns on the TransposeOp so
 /// that they can be picked up by the Canonicalization framework.
 void TransposeOp::getCanonicalizationPatterns(OwningRewritePatternList &results,
                                               MLIRContext *context) {
   results.insert<SimplifyRedundantTranspose>(context);

mlir/examples/toy/Ch6/mlir/LowerToAffineLoops.cpp

 // ToyToAffine RewritePatterns: Binary operations
 //===----------------------------------------------------------------------===//
 
 template <typename BinaryOp, typename LoweredBinaryOp>
 struct BinaryOpLowering : public ConversionPattern {
   BinaryOpLowering(MLIRContext *ctx)
       : ConversionPattern(BinaryOp::getOperationName(), 1, ctx) {}
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const final {
     auto loc = op->getLoc();
     lowerOpToLoops(
         op, operands, rewriter,
         [loc](PatternRewriter &rewriter, ArrayRef<Value> memRefOperands,
               ArrayRef<Value> loopIvs) {
           // Generate an adaptor for the remapped operands of the BinaryOp. This
           // allows for using the nice named accessors that are generated by the
           // ODS.
           typename BinaryOp::OperandAdaptor binaryAdaptor(memRefOperands);
 
           // Generate loads for the element of 'lhs' and 'rhs' at the inner
           // loop.
           auto loadedLhs =
               rewriter.create<AffineLoadOp>(loc, binaryAdaptor.lhs(), loopIvs);
           auto loadedRhs =
               rewriter.create<AffineLoadOp>(loc, binaryAdaptor.rhs(), loopIvs);
 
           // Create the binary operation performed on the loaded values.
           return rewriter.create<LoweredBinaryOp>(loc, loadedLhs, loadedRhs);
         });
-    return matchSuccess();
+    return success();
   }
 };
 using AddOpLowering = BinaryOpLowering<toy::AddOp, AddFOp>;
 using MulOpLowering = BinaryOpLowering<toy::MulOp, MulFOp>;
 
 //===----------------------------------------------------------------------===//
 // ToyToAffine RewritePatterns: Constant operations
 //===----------------------------------------------------------------------===//
 
 struct ConstantOpLowering : public OpRewritePattern<toy::ConstantOp> {
   using OpRewritePattern<toy::ConstantOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(toy::ConstantOp op,
-                                     PatternRewriter &rewriter) const final {
+  LogicalResult matchAndRewrite(toy::ConstantOp op,
     DenseElementsAttr constantValue = op.value();
     Location loc = op.getLoc();
 
     // When lowering the constant operation, we allocate and assign the constant
     // values to a corresponding memref allocation.
     auto tensorType = op.getType().cast<TensorType>();
     auto memRefType = convertTensorToMemRef(tensorType);
     auto alloc = insertAllocAndDealloc(memRefType, loc, rewriter);
std::function<void(uint64_t)> storeElements = [&](uint64_t dimension) {
       }
     };
 
     // Start the element storing recursion from the first dimension.
     storeElements(/*dimension=*/0);
 
     // Replace this operation with the generated alloc.
     rewriter.replaceOp(op, alloc);
-    return matchSuccess();
+    return success();
   }
 };
 
 //===----------------------------------------------------------------------===//
 // ToyToAffine RewritePatterns: Return operations
 //===----------------------------------------------------------------------===//
 
 struct ReturnOpLowering : public OpRewritePattern<toy::ReturnOp> {
   using OpRewritePattern<toy::ReturnOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(toy::ReturnOp op,
-                                     PatternRewriter &rewriter) const final {
+  LogicalResult matchAndRewrite(toy::ReturnOp op,
     // During this lowering, we expect that all function calls have been
     // inlined.
     if (op.hasOperand())
-      return matchFailure();
+      return failure();
 
     // We lower "toy.return" directly to "std.return".
     rewriter.replaceOpWithNewOp<ReturnOp>(op);
-    return matchSuccess();
+    return success();
   }
 };
 
 //===----------------------------------------------------------------------===//
 // ToyToAffine RewritePatterns: Transpose operations
 //===----------------------------------------------------------------------===//
 
 struct TransposeOpLowering : public ConversionPattern {
   TransposeOpLowering(MLIRContext *ctx)
       : ConversionPattern(toy::TransposeOp::getOperationName(), 1, ctx) {}
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const final {
     auto loc = op->getLoc();
     lowerOpToLoops(
         op, operands, rewriter,
         [loc](PatternRewriter &rewriter, ArrayRef<Value> memRefOperands,
               ArrayRef<Value> loopIvs) {
           // Generate an adaptor for the remapped operands of the TransposeOp.
           // This allows for using the nice named accessors that are generated
           // by the ODS.
           toy::TransposeOpOperandAdaptor transposeAdaptor(memRefOperands);
           Value input = transposeAdaptor.input();
 
           // Transpose the elements by generating a load from the reverse
           // indices.
           SmallVector<Value, 2> reverseIvs(llvm::reverse(loopIvs));
           return rewriter.create<AffineLoadOp>(loc, input, reverseIvs);
         });
-    return matchSuccess();
+    return success();
   }
 };
 
 } // end anonymous namespace.
 
 //===----------------------------------------------------------------------===//
 // ToyToAffineLoweringPass
 //===----------------------------------------------------------------------===//

mlir/examples/toy/Ch6/mlir/LowerToLLVM.cpp

 namespace {
 /// Lowers `toy.print` to a loop nest calling `printf` on each of the individual
 /// elements of the array.
 class PrintOpLowering : public ConversionPattern {
 public:
   explicit PrintOpLowering(MLIRContext *context)
       : ConversionPattern(toy::PrintOp::getOperationName(), 1, context) {}
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto memRefType = (*op->operand_type_begin()).cast<MemRefType>();
     auto memRefShape = memRefType.getShape();
     auto loc = op->getLoc();
     auto *llvmDialect =
         op->getContext()->getRegisteredDialect<LLVM::LLVMDialect>();
     assert(llvmDialect && "expected llvm dialect to be registered");
matchAndRewrite(Operation *op, ArrayRef<Value> operands,
     // Generate a call to printf for the current element of the loop.
     auto printOp = cast<toy::PrintOp>(op);
     auto elementLoad = rewriter.create<LoadOp>(loc, printOp.input(), loopIvs);
     rewriter.create<CallOp>(loc, printfRef, rewriter.getIntegerType(32),
                             ArrayRef<Value>({formatSpecifierCst, elementLoad}));
 
     // Notify the rewriter that this operation has been removed.
     rewriter.eraseOp(op);
-    return matchSuccess();
+    return success();
   }
 
 private:
   /// Return a symbol reference to the printf function, inserting it into the
   /// module if necessary.
   static FlatSymbolRefAttr getOrInsertPrintf(PatternRewriter &rewriter,
                                              ModuleOp module,
                                              LLVM::LLVMDialect *llvmDialect) {

mlir/examples/toy/Ch6/mlir/ToyCombine.cpp

struct SimplifyRedundantTranspose : public mlir::OpRewritePattern<TransposeOp> {
   /// The "benefit" is used by the framework to order the patterns and process
   /// them in order of profitability.
   SimplifyRedundantTranspose(mlir::MLIRContext *context)
       : OpRewritePattern<TransposeOp>(context, /*benefit=*/1) {}
 
   /// This method attempts to match a pattern and rewrite it. The rewriter
   /// argument is the orchestrator of the sequence of rewrites. The pattern is
   /// expected to interact with it to perform any changes to the IR from here.
-  mlir::PatternMatchResult
+  mlir::LogicalResult
   matchAndRewrite(TransposeOp op,
                   mlir::PatternRewriter &rewriter) const override {
     // Look through the input of the current transpose.
     mlir::Value transposeInput = op.getOperand();
     TransposeOp transposeInputOp =
         llvm::dyn_cast_or_null<TransposeOp>(transposeInput.getDefiningOp());
 
     // Input defined by another transpose? If not, no match.
     if (!transposeInputOp)
-      return matchFailure();
+      return failure();
 
     // Otherwise, we have a redundant transpose. Use the rewriter.
     rewriter.replaceOp(op, {transposeInputOp.getOperand()});
-    return matchSuccess();
+    return success();
   }
 };
 
 /// Register our patterns as "canonicalization" patterns on the TransposeOp so
 /// that they can be picked up by the Canonicalization framework.
 void TransposeOp::getCanonicalizationPatterns(OwningRewritePatternList &results,
                                               MLIRContext *context) {
   results.insert<SimplifyRedundantTranspose>(context);

mlir/examples/toy/Ch7/mlir/LowerToAffineLoops.cpp

 // ToyToAffine RewritePatterns: Binary operations
 //===----------------------------------------------------------------------===//
 
 template <typename BinaryOp, typename LoweredBinaryOp>
 struct BinaryOpLowering : public ConversionPattern {
   BinaryOpLowering(MLIRContext *ctx)
       : ConversionPattern(BinaryOp::getOperationName(), 1, ctx) {}
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const final {
     auto loc = op->getLoc();
     lowerOpToLoops(
         op, operands, rewriter,
         [loc](PatternRewriter &rewriter, ArrayRef<Value> memRefOperands,
               ArrayRef<Value> loopIvs) {
           // Generate an adaptor for the remapped operands of the BinaryOp. This
           // allows for using the nice named accessors that are generated by the
           // ODS.
           typename BinaryOp::OperandAdaptor binaryAdaptor(memRefOperands);
 
           // Generate loads for the element of 'lhs' and 'rhs' at the inner
           // loop.
           auto loadedLhs =
               rewriter.create<AffineLoadOp>(loc, binaryAdaptor.lhs(), loopIvs);
           auto loadedRhs =
               rewriter.create<AffineLoadOp>(loc, binaryAdaptor.rhs(), loopIvs);
 
           // Create the binary operation performed on the loaded values.
           return rewriter.create<LoweredBinaryOp>(loc, loadedLhs, loadedRhs);
         });
-    return matchSuccess();
+    return success();
   }
 };
 using AddOpLowering = BinaryOpLowering<toy::AddOp, AddFOp>;
 using MulOpLowering = BinaryOpLowering<toy::MulOp, MulFOp>;
 
 //===----------------------------------------------------------------------===//
 // ToyToAffine RewritePatterns: Constant operations
 //===----------------------------------------------------------------------===//
 
 struct ConstantOpLowering : public OpRewritePattern<toy::ConstantOp> {
   using OpRewritePattern<toy::ConstantOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(toy::ConstantOp op,
-                                     PatternRewriter &rewriter) const final {
+  LogicalResult matchAndRewrite(toy::ConstantOp op,
     DenseElementsAttr constantValue = op.value();
     Location loc = op.getLoc();
 
     // When lowering the constant operation, we allocate and assign the constant
     // values to a corresponding memref allocation.
     auto tensorType = op.getType().cast<TensorType>();
     auto memRefType = convertTensorToMemRef(tensorType);
     auto alloc = insertAllocAndDealloc(memRefType, loc, rewriter);
std::function<void(uint64_t)> storeElements = [&](uint64_t dimension) {
       }
     };
 
     // Start the element storing recursion from the first dimension.
     storeElements(/*dimension=*/0);
 
     // Replace this operation with the generated alloc.
     rewriter.replaceOp(op, alloc);
-    return matchSuccess();
+    return success();
   }
 };
 
 //===----------------------------------------------------------------------===//
 // ToyToAffine RewritePatterns: Return operations
 //===----------------------------------------------------------------------===//
 
 struct ReturnOpLowering : public OpRewritePattern<toy::ReturnOp> {
   using OpRewritePattern<toy::ReturnOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(toy::ReturnOp op,
-                                     PatternRewriter &rewriter) const final {
+  LogicalResult matchAndRewrite(toy::ReturnOp op,
     // During this lowering, we expect that all function calls have been
     // inlined.
     if (op.hasOperand())
-      return matchFailure();
+      return failure();
 
     // We lower "toy.return" directly to "std.return".
     rewriter.replaceOpWithNewOp<ReturnOp>(op);
-    return matchSuccess();
+    return success();
   }
 };
 
 //===----------------------------------------------------------------------===//
 // ToyToAffine RewritePatterns: Transpose operations
 //===----------------------------------------------------------------------===//
 
 struct TransposeOpLowering : public ConversionPattern {
   TransposeOpLowering(MLIRContext *ctx)
       : ConversionPattern(toy::TransposeOp::getOperationName(), 1, ctx) {}
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const final {
     auto loc = op->getLoc();
     lowerOpToLoops(
         op, operands, rewriter,
         [loc](PatternRewriter &rewriter, ArrayRef<Value> memRefOperands,
               ArrayRef<Value> loopIvs) {
           // Generate an adaptor for the remapped operands of the TransposeOp.
           // This allows for using the nice named accessors that are generated
           // by the ODS.
           toy::TransposeOpOperandAdaptor transposeAdaptor(memRefOperands);
           Value input = transposeAdaptor.input();
 
           // Transpose the elements by generating a load from the reverse
           // indices.
           SmallVector<Value, 2> reverseIvs(llvm::reverse(loopIvs));
           return rewriter.create<AffineLoadOp>(loc, input, reverseIvs);
         });
-    return matchSuccess();
+    return success();
   }
 };
 
 } // end anonymous namespace.
 
 //===----------------------------------------------------------------------===//
 // ToyToAffineLoweringPass
 //===----------------------------------------------------------------------===//

mlir/examples/toy/Ch7/mlir/LowerToLLVM.cpp

 namespace {
 /// Lowers `toy.print` to a loop nest calling `printf` on each of the individual
 /// elements of the array.
 class PrintOpLowering : public ConversionPattern {
 public:
   explicit PrintOpLowering(MLIRContext *context)
       : ConversionPattern(toy::PrintOp::getOperationName(), 1, context) {}
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto memRefType = (*op->operand_type_begin()).cast<MemRefType>();
     auto memRefShape = memRefType.getShape();
     auto loc = op->getLoc();
     auto *llvmDialect =
         op->getContext()->getRegisteredDialect<LLVM::LLVMDialect>();
     assert(llvmDialect && "expected llvm dialect to be registered");
matchAndRewrite(Operation *op, ArrayRef<Value> operands,
     // Generate a call to printf for the current element of the loop.
     auto printOp = cast<toy::PrintOp>(op);
     auto elementLoad = rewriter.create<LoadOp>(loc, printOp.input(), loopIvs);
     rewriter.create<CallOp>(loc, printfRef, rewriter.getIntegerType(32),
                             ArrayRef<Value>({formatSpecifierCst, elementLoad}));
 
     // Notify the rewriter that this operation has been removed.
     rewriter.eraseOp(op);
-    return matchSuccess();
+    return success();
   }
 
 private:
   /// Return a symbol reference to the printf function, inserting it into the
   /// module if necessary.
   static FlatSymbolRefAttr getOrInsertPrintf(PatternRewriter &rewriter,
                                              ModuleOp module,
                                              LLVM::LLVMDialect *llvmDialect) {

mlir/examples/toy/Ch7/mlir/ToyCombine.cpp

struct SimplifyRedundantTranspose : public mlir::OpRewritePattern<TransposeOp> {
   /// The "benefit" is used by the framework to order the patterns and process
   /// them in order of profitability.
   SimplifyRedundantTranspose(mlir::MLIRContext *context)
       : OpRewritePattern<TransposeOp>(context, /*benefit=*/1) {}
 
   /// This method attempts to match a pattern and rewrite it. The rewriter
   /// argument is the orchestrator of the sequence of rewrites. The pattern is
   /// expected to interact with it to perform any changes to the IR from here.
-  mlir::PatternMatchResult
+  mlir::LogicalResult
   matchAndRewrite(TransposeOp op,
                   mlir::PatternRewriter &rewriter) const override {
     // Look through the input of the current transpose.
     mlir::Value transposeInput = op.getOperand();
     TransposeOp transposeInputOp =
         llvm::dyn_cast_or_null<TransposeOp>(transposeInput.getDefiningOp());
 
     // Input defined by another transpose? If not, no match.
     if (!transposeInputOp)
-      return matchFailure();
+      return failure();
 
     // Otherwise, we have a redundant transpose. Use the rewriter.
     rewriter.replaceOp(op, {transposeInputOp.getOperand()});
-    return matchSuccess();
+    return success();
   }
 };
 
 /// Register our patterns as "canonicalization" patterns on the TransposeOp so
 /// that they can be picked up by the Canonicalization framework.
 void TransposeOp::getCanonicalizationPatterns(OwningRewritePatternList &results,
                                               MLIRContext *context) {
   results.insert<SimplifyRedundantTranspose>(context);

mlir/include/mlir/Dialect/Linalg/Transforms/LinalgTransformPatterns.td

 // In the future, tile sizes should be derived from op properties + machine
 // description but we do not need to wait on this to start having useful
 // patterns.
 class TileAndFuseLinalgOp<
     list<int> sizes, list<int> operandIndices, string value> : NativeCodeCall<
   "if (failed(tileAndFuseLinalgOpAndSetMarker($_builder, op, {" #
   StrJoinInt<sizes>.result # "}, {" # StrJoinInt<operandIndices>.result # "}," #
       " \"" # value # "\")))" #
-  "  return matchFailure();">;
+  "  return failure();">;
 
 //===----------------------------------------------------------------------===//
 // Linalg tiling patterns.
 //===----------------------------------------------------------------------===//
 //
 // In the future, tile sizes should be derived from op properties + machine
 // description but we do not need to wait on this to start having useful
 // patterns.
 // `permutation` is an optional parameter to specify the ordering of the
 // tiled loops. If provided, it must be a list of integers with the same number
 // of elements as `sizes`.
 class TileLinalgOp<list<int> sizes, string value, list<int> permutation=[]> :
   NativeCodeCall<
     "if (failed(tileLinalgOpAndSetMarker($_builder, op, {" #
     StrJoinInt<sizes>.result # "}, \"" # value # "\", {" #
     StrJoinInt<permutation>.result # "})))" #
-    "  return matchFailure();">;
+    "  return failure();">;
 
 //===----------------------------------------------------------------------===//
 // Linalg to loop patterns.
 //===----------------------------------------------------------------------===//
 class LinalgOpToLoops<string OpType> : NativeCodeCall<
   "if (failed(linalgOpToLoops<" # OpType # ">($_builder, op))) " #
-  "  return matchFailure();">;
+  "  return failure();">;
 
 class LinalgOpToParallelLoops<string OpType> : NativeCodeCall<
   "if (failed(linalgOpToParallelLoops<" # OpType # ">($_builder, op))) " #
-  "  return matchFailure();">;
+  "  return failure();">;
 
 class LinalgOpToAffineLoops<string OpType> : NativeCodeCall<
   "if (failed(linalgOpToAffineLoops<" # OpType # ">($_builder, op))) " #
-  "  return matchFailure();">;
+  "  return failure();">;
 
 //===----------------------------------------------------------------------===//
 // Linalg to vector patterns precondition and DRR.
 //===----------------------------------------------------------------------===//
 def PreconditionVectorizeLinalgOp : CPred<
   "succeeded(vectorizeLinalgOpPrecondition(op))">;
 def VectorizeLinalgOp : NativeCodeCall<
   "vectorizeLinalgOp($_builder, op)">;

mlir/include/mlir/IR/PatternMatch.h

bool operator<(const PatternBenefit &rhs) const {
     return representation < rhs.representation;
   }
 
 private:
   PatternBenefit() : representation(ImpossibleToMatchSentinel) {}
   unsigned short representation;
 };
 
-/// This is the type returned by a pattern match.
-/// TODO: Replace usages with LogicalResult directly.
-using PatternMatchResult = LogicalResult;
-
 //===----------------------------------------------------------------------===//
 // Pattern class
 //===----------------------------------------------------------------------===//
 
 /// Instances of Pattern can be matched against SSA IR.  These matches get used
 /// in ways dependent on their subclasses and the driver doing the matching.
 /// For example, RewritePatterns implement a rewrite from one matched pattern
 /// to a replacement DAG tile.
public:
   OperationName getRootKind() const { return rootKind; }
 
   //===--------------------------------------------------------------------===//
   // Implementation hooks for patterns to implement.
   //===--------------------------------------------------------------------===//
 
   /// Attempt to match against code rooted at the specified operation,
   /// which is the same operation code as getRootKind().
-  virtual PatternMatchResult match(Operation *op) const = 0;
+  virtual LogicalResult match(Operation *op) const = 0;
 
   virtual ~Pattern() {}
 
-  //===--------------------------------------------------------------------===//
-  // Helper methods to simplify pattern implementations
-  //===--------------------------------------------------------------------===//
-
-  /// Return a result, indicating that no match was found.
-  PatternMatchResult matchFailure() const { return failure(); }
-
-  /// This method indicates that a match was found.
-  PatternMatchResult matchSuccess() const { return success(); }
-
 protected:
   /// Patterns must specify the root operation name they match against, and can
   /// also specify the benefit of the pattern matching.
   Pattern(StringRef rootName, PatternBenefit benefit, MLIRContext *context);
 
 private:
   const OperationName rootKind;
   const PatternBenefit benefit;
public:
   /// Rewrite the IR rooted at the specified operation with the result of
   /// this pattern, generating any new operations with the specified
   /// builder.  If an unexpected error is encountered (an internal
   /// compiler error), it is emitted through the normal MLIR diagnostic
   /// hooks and the IR is left in a valid state.
   virtual void rewrite(Operation *op, PatternRewriter &rewriter) const;
 
   /// Attempt to match against code rooted at the specified operation,
-  /// which is the same operation code as getRootKind().  On failure, this
-  /// returns a None value.  On success, it returns a (possibly null)
-  /// pattern-specific state wrapped in an Optional.  This state is passed back
-  /// into the rewrite function if this match is selected.
-  PatternMatchResult match(Operation *op) const override;
+  /// which is the same operation code as getRootKind().
+  LogicalResult match(Operation *op) const override;
 
   /// Attempt to match against code rooted at the specified operation,
   /// which is the same operation code as getRootKind(). If successful, this
   /// function will automatically perform the rewrite.
-  virtual PatternMatchResult matchAndRewrite(Operation *op,
-                                             PatternRewriter &rewriter) const {
+  virtual LogicalResult matchAndRewrite(Operation *op,
     if (succeeded(match(op))) {
       rewrite(op, rewriter);
-      return matchSuccess();
+      return success();
     }
-    return matchFailure();
+    return failure();
   }
 
   /// Return a list of operations that may be generated when rewriting an
   /// operation instance with this pattern.
   ArrayRef<OperationName> getGeneratedOps() const { return generatedOps; }
 
 protected:
   /// Patterns must specify the root operation name they match against, and can
template <typename SourceOp> struct OpRewritePattern : public RewritePattern {
   /// also specify the benefit of the pattern matching.
   OpRewritePattern(MLIRContext *context, PatternBenefit benefit = 1)
       : RewritePattern(SourceOp::getOperationName(), benefit, context) {}
 
   /// Wrappers around the RewritePattern methods that pass the derived op type.
   void rewrite(Operation *op, PatternRewriter &rewriter) const final {
     rewrite(cast<SourceOp>(op), rewriter);
   }
-  PatternMatchResult match(Operation *op) const final {
+  LogicalResult match(Operation *op) const final {
     return match(cast<SourceOp>(op));
   }
-  PatternMatchResult matchAndRewrite(Operation *op,
-                                     PatternRewriter &rewriter) const final {
+  LogicalResult matchAndRewrite(Operation *op,
     return matchAndRewrite(cast<SourceOp>(op), rewriter);
   }
 
   /// Rewrite and Match methods that operate on the SourceOp type. These must be
   /// overridden by the derived pattern class.
   virtual void rewrite(SourceOp op, PatternRewriter &rewriter) const {
     llvm_unreachable("must override rewrite or matchAndRewrite");
   }
-  virtual PatternMatchResult match(SourceOp op) const {
+  virtual LogicalResult match(SourceOp op) const {
     llvm_unreachable("must override match or matchAndRewrite");
   }
-  virtual PatternMatchResult matchAndRewrite(SourceOp op,
-                                             PatternRewriter &rewriter) const {
+  virtual LogicalResult matchAndRewrite(SourceOp op,
     if (succeeded(match(op))) {
       rewrite(op, rewriter);
-      return matchSuccess();
+      return success();
     }
-    return matchFailure();
+    return failure();
   }
 };
 
 //===----------------------------------------------------------------------===//
 // PatternRewriter class
 //===----------------------------------------------------------------------===//
 
 /// This class coordinates the application of a pattern to the current function,

mlir/include/mlir/Transforms/DialectConversion.h

public:
   /// operations. This function should not fail. If some specific cases of
   /// the operation are not supported, these cases should not be matched.
   virtual void rewrite(Operation *op, ArrayRef<Value> operands,
                        ConversionPatternRewriter &rewriter) const {
     llvm_unreachable("unimplemented rewrite");
   }
 
   /// Hook for derived classes to implement combined matching and rewriting.
-  virtual PatternMatchResult
+  virtual LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const {
     if (failed(match(op)))
-      return matchFailure();
+      return failure();
     rewrite(op, operands, rewriter);
-    return matchSuccess();
+    return success();
   }
 
   /// Attempt to match and rewrite the IR root at the specified operation.
-  PatternMatchResult matchAndRewrite(Operation *op,
-                                     PatternRewriter &rewriter) const final;
+  LogicalResult matchAndRewrite(Operation *op,
 
 private:
   using RewritePattern::rewrite;
 };
 
 /// OpConversionPattern is a wrapper around ConversionPattern that allows for
 /// matching and rewriting against an instance of a derived operation class as
 /// opposed to a raw Operation.
 template <typename SourceOp>
 struct OpConversionPattern : public ConversionPattern {
   OpConversionPattern(MLIRContext *context, PatternBenefit benefit = 1)
       : ConversionPattern(SourceOp::getOperationName(), benefit, context) {}
 
   /// Wrappers around the ConversionPattern methods that pass the derived op
   /// type.
   void rewrite(Operation *op, ArrayRef<Value> operands,
                ConversionPatternRewriter &rewriter) const final {
     rewrite(cast<SourceOp>(op), operands, rewriter);
   }
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const final {
     return matchAndRewrite(cast<SourceOp>(op), operands, rewriter);
   }
 
   // TODO(b/142763075): Use OperandAdaptor when it supports access to unnamed
   // operands.
 
   /// Rewrite and Match methods that operate on the SourceOp type. These must be
   /// overridden by the derived pattern class.
   virtual void rewrite(SourceOp op, ArrayRef<Value> operands,
                        ConversionPatternRewriter &rewriter) const {
     llvm_unreachable("must override matchAndRewrite or a rewrite method");
   }
 
-  virtual PatternMatchResult
+  virtual LogicalResult
   matchAndRewrite(SourceOp op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const {
     if (failed(match(op)))
-      return matchFailure();
+      return failure();
     rewrite(op, operands, rewriter);
-    return matchSuccess();
+    return success();
   }
 
 private:
   using ConversionPattern::matchAndRewrite;
 };
 
 /// Add a pattern to the given pattern list to convert the signature of a FuncOp
 /// with the given type converter.

mlir/lib/Conversion/AffineToStandard/AffineToStandard.cpp

return lowerAffineMapMax(builder, op.getLoc(), op.getLowerBoundMap(),
                            op.getLowerBoundOperands());
 }
 
 namespace {
 class AffineMinLowering : public OpRewritePattern<AffineMinOp> {
 public:
   using OpRewritePattern<AffineMinOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(AffineMinOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(AffineMinOp op,
     Value reduced =
         lowerAffineMapMin(rewriter, op.getLoc(), op.map(), op.operands());
     if (!reduced)
-      return matchFailure();
+      return failure();
 
     rewriter.replaceOp(op, reduced);
-    return matchSuccess();
+    return success();
   }
 };
 
 class AffineMaxLowering : public OpRewritePattern<AffineMaxOp> {
 public:
   using OpRewritePattern<AffineMaxOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(AffineMaxOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(AffineMaxOp op,
     Value reduced =
         lowerAffineMapMax(rewriter, op.getLoc(), op.map(), op.operands());
     if (!reduced)
-      return matchFailure();
+      return failure();
 
     rewriter.replaceOp(op, reduced);
-    return matchSuccess();
+    return success();
   }
 };
 
 /// Affine terminators are removed.
 class AffineTerminatorLowering : public OpRewritePattern<AffineTerminatorOp> {
 public:
   using OpRewritePattern<AffineTerminatorOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(AffineTerminatorOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(AffineTerminatorOp op,
     rewriter.replaceOpWithNewOp<loop::YieldOp>(op);
-    return matchSuccess();
+    return success();
   }
 };
 
 class AffineForLowering : public OpRewritePattern<AffineForOp> {
 public:
   using OpRewritePattern<AffineForOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(AffineForOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(AffineForOp op,
     Location loc = op.getLoc();
     Value lowerBound = lowerAffineLowerBound(op, rewriter);
     Value upperBound = lowerAffineUpperBound(op, rewriter);
     Value step = rewriter.create<ConstantIndexOp>(loc, op.getStep());
     auto f = rewriter.create<loop::ForOp>(loc, lowerBound, upperBound, step);
     f.region().getBlocks().clear();
     rewriter.inlineRegionBefore(op.region(), f.region(), f.region().end());
     rewriter.eraseOp(op);
-    return matchSuccess();
+    return success();
   }
 };
 
 class AffineIfLowering : public OpRewritePattern<AffineIfOp> {
 public:
   using OpRewritePattern<AffineIfOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(AffineIfOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(AffineIfOp op,
     auto loc = op.getLoc();
 
     // Now we just have to handle the condition logic.
     auto integerSet = op.getIntegerSet();
     Value zeroConstant = rewriter.create<ConstantIndexOp>(loc, 0);
     SmallVector<Value, 8> operands(op.getOperands());
     auto operandsRef = llvm::makeArrayRef(operands);
 
     // Calculate cond as a conjunction without short-circuiting.
     Value cond = nullptr;
     for (unsigned i = 0, e = integerSet.getNumConstraints(); i < e; ++i) {
       AffineExpr constraintExpr = integerSet.getConstraint(i);
       bool isEquality = integerSet.isEq(i);
 
       // Build and apply an affine expression
       auto numDims = integerSet.getNumDims();
       Value affResult = expandAffineExpr(rewriter, loc, constraintExpr,
                                          operandsRef.take_front(numDims),
                                          operandsRef.drop_front(numDims));
       if (!affResult)
-        return matchFailure();
+        return failure();
       auto pred = isEquality ? CmpIPredicate::eq : CmpIPredicate::sge;
       Value cmpVal =
           rewriter.create<CmpIOp>(loc, pred, affResult, zeroConstant);
       cond =
           cond ? rewriter.create<AndOp>(loc, cond, cmpVal).getResult() : cmpVal;
     }
     cond = cond ? cond
                 : rewriter.create<ConstantIntOp>(loc, /*value=*/1, /*width=*/1);
 
     bool hasElseRegion = !op.elseRegion().empty();
     auto ifOp = rewriter.create<loop::IfOp>(loc, cond, hasElseRegion);
     rewriter.inlineRegionBefore(op.thenRegion(), &ifOp.thenRegion().back());
     ifOp.thenRegion().back().erase();
     if (hasElseRegion) {
       rewriter.inlineRegionBefore(op.elseRegion(), &ifOp.elseRegion().back());
       ifOp.elseRegion().back().erase();
     }
 
     // Ok, we're done!
     rewriter.eraseOp(op);
-    return matchSuccess();
+    return success();
   }
 };
 
 /// Convert an "affine.apply" operation into a sequence of arithmetic
 /// operations using the StandardOps dialect.
 class AffineApplyLowering : public OpRewritePattern<AffineApplyOp> {
 public:
   using OpRewritePattern<AffineApplyOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(AffineApplyOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(AffineApplyOp op,
     auto maybeExpandedMap =
         expandAffineMap(rewriter, op.getLoc(), op.getAffineMap(),
                         llvm::to_vector<8>(op.getOperands()));
     if (!maybeExpandedMap)
-      return matchFailure();
+      return failure();
     rewriter.replaceOp(op, *maybeExpandedMap);
-    return matchSuccess();
+    return success();
   }
 };
 
 /// Apply the affine map from an 'affine.load' operation to its operands, and
 /// feed the results to a newly created 'std.load' operation (which replaces the
 /// original 'affine.load').
 class AffineLoadLowering : public OpRewritePattern<AffineLoadOp> {
 public:
   using OpRewritePattern<AffineLoadOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(AffineLoadOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(AffineLoadOp op,
     // Expand affine map from 'affineLoadOp'.
     SmallVector<Value, 8> indices(op.getMapOperands());
     auto resultOperands =
         expandAffineMap(rewriter, op.getLoc(), op.getAffineMap(), indices);
     if (!resultOperands)
-      return matchFailure();
+      return failure();
 
     // Build std.load memref[expandedMap.results].
     rewriter.replaceOpWithNewOp<LoadOp>(op, op.getMemRef(), *resultOperands);
-    return matchSuccess();
+    return success();
   }
 };
 
 /// Apply the affine map from an 'affine.prefetch' operation to its operands,
 /// and feed the results to a newly created 'std.prefetch' operation (which
 /// replaces the original 'affine.prefetch').
 class AffinePrefetchLowering : public OpRewritePattern<AffinePrefetchOp> {
 public:
   using OpRewritePattern<AffinePrefetchOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(AffinePrefetchOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(AffinePrefetchOp op,
     // Expand affine map from 'affinePrefetchOp'.
     SmallVector<Value, 8> indices(op.getMapOperands());
     auto resultOperands =
         expandAffineMap(rewriter, op.getLoc(), op.getAffineMap(), indices);
     if (!resultOperands)
-      return matchFailure();
+      return failure();
 
     // Build std.prefetch memref[expandedMap.results].
     rewriter.replaceOpWithNewOp<PrefetchOp>(
         op, op.memref(), *resultOperands, op.isWrite(),
         op.localityHint().getZExtValue(), op.isDataCache());
-    return matchSuccess();
+    return success();
   }
 };
 
 /// Apply the affine map from an 'affine.store' operation to its operands, and
 /// feed the results to a newly created 'std.store' operation (which replaces
 /// the original 'affine.store').
 class AffineStoreLowering : public OpRewritePattern<AffineStoreOp> {
 public:
   using OpRewritePattern<AffineStoreOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(AffineStoreOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(AffineStoreOp op,
     // Expand affine map from 'affineStoreOp'.
     SmallVector<Value, 8> indices(op.getMapOperands());
     auto maybeExpandedMap =
         expandAffineMap(rewriter, op.getLoc(), op.getAffineMap(), indices);
     if (!maybeExpandedMap)
-      return matchFailure();
+      return failure();
 
     // Build std.store valueToStore, memref[expandedMap.results].
     rewriter.replaceOpWithNewOp<StoreOp>(op, op.getValueToStore(),
                                          op.getMemRef(), *maybeExpandedMap);
-    return matchSuccess();
+    return success();
   }
 };
 
 /// Apply the affine maps from an 'affine.dma_start' operation to each of their
 /// respective map operands, and feed the results to a newly created
 /// 'std.dma_start' operation (which replaces the original 'affine.dma_start').
 class AffineDmaStartLowering : public OpRewritePattern<AffineDmaStartOp> {
 public:
   using OpRewritePattern<AffineDmaStartOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(AffineDmaStartOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(AffineDmaStartOp op,
     SmallVector<Value, 8> operands(op.getOperands());
     auto operandsRef = llvm::makeArrayRef(operands);
 
     // Expand affine map for DMA source memref.
     auto maybeExpandedSrcMap = expandAffineMap(
         rewriter, op.getLoc(), op.getSrcMap(),
         operandsRef.drop_front(op.getSrcMemRefOperandIndex() + 1));
     if (!maybeExpandedSrcMap)
-      return matchFailure();
+      return failure();
     // Expand affine map for DMA destination memref.
     auto maybeExpandedDstMap = expandAffineMap(
         rewriter, op.getLoc(), op.getDstMap(),
         operandsRef.drop_front(op.getDstMemRefOperandIndex() + 1));
     if (!maybeExpandedDstMap)
-      return matchFailure();
+      return failure();
     // Expand affine map for DMA tag memref.
     auto maybeExpandedTagMap = expandAffineMap(
         rewriter, op.getLoc(), op.getTagMap(),
         operandsRef.drop_front(op.getTagMemRefOperandIndex() + 1));
     if (!maybeExpandedTagMap)
-      return matchFailure();
+      return failure();
 
     // Build std.dma_start operation with affine map results.
     rewriter.replaceOpWithNewOp<DmaStartOp>(
         op, op.getSrcMemRef(), *maybeExpandedSrcMap, op.getDstMemRef(),
         *maybeExpandedDstMap, op.getNumElements(), op.getTagMemRef(),
         *maybeExpandedTagMap, op.getStride(), op.getNumElementsPerStride());
-    return matchSuccess();
+    return success();
   }
 };
 
 /// Apply the affine map from an 'affine.dma_wait' operation tag memref,
 /// and feed the results to a newly created 'std.dma_wait' operation (which
 /// replaces the original 'affine.dma_wait').
 class AffineDmaWaitLowering : public OpRewritePattern<AffineDmaWaitOp> {
 public:
   using OpRewritePattern<AffineDmaWaitOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(AffineDmaWaitOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(AffineDmaWaitOp op,
     // Expand affine map for DMA tag memref.
     SmallVector<Value, 8> indices(op.getTagIndices());
     auto maybeExpandedTagMap =
         expandAffineMap(rewriter, op.getLoc(), op.getTagMap(), indices);
     if (!maybeExpandedTagMap)
-      return matchFailure();
+      return failure();
 
     // Build std.dma_wait operation with affine map results.
     rewriter.replaceOpWithNewOp<DmaWaitOp>(
         op, op.getTagMemRef(), *maybeExpandedTagMap, op.getNumElements());
-    return matchSuccess();
+    return success();
   }
 };
 
 } // end namespace
 
 void mlir::populateAffineToStdConversionPatterns(
     OwningRewritePatternList &patterns, MLIRContext *ctx) {
   // clang-format off

mlir/lib/Conversion/GPUCommon/IndexIntrinsicsOpLowering.h

 
 public:
   explicit GPUIndexIntrinsicOpLowering(LLVMTypeConverter &lowering_)
       : ConvertToLLVMPattern(Op::getOperationName(),
                              lowering_.getDialect()->getContext(), lowering_),
         indexBitwidth(getIndexBitWidth(lowering_)) {}
 
   // Convert the kernel arguments to an LLVM type, preserve the rest.
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto loc = op->getLoc();
     auto dialect = typeConverter.getDialect();
     Value newOp;
     switch (dimensionToIndex(cast<Op>(op))) {
     case X:
       newOp = rewriter.create<XOp>(loc, LLVM::LLVMType::getInt32Ty(dialect));
       break;
     case Y:
       newOp = rewriter.create<YOp>(loc, LLVM::LLVMType::getInt32Ty(dialect));
       break;
     case Z:
       newOp = rewriter.create<ZOp>(loc, LLVM::LLVMType::getInt32Ty(dialect));
       break;
     default:
-      return matchFailure();
+      return failure();
     }
 
     if (indexBitwidth > 32) {
       newOp = rewriter.create<LLVM::SExtOp>(
           loc, LLVM::LLVMType::getIntNTy(dialect, indexBitwidth), newOp);
     } else if (indexBitwidth < 32) {
       newOp = rewriter.create<LLVM::TruncOp>(
           loc, LLVM::LLVMType::getIntNTy(dialect, indexBitwidth), newOp);
     }
 
     rewriter.replaceOp(op, {newOp});
-    return matchSuccess();
+    return success();
   }
 };
 
 } // namespace mlir
 
 #endif // MLIR_CONVERSION_GPUCOMMON_INDEXINTRINSICSOPLOWERING_H_

mlir/lib/Conversion/GPUCommon/OpToFuncCallLowering.h

 struct OpToFuncCallLowering : public ConvertToLLVMPattern {
 public:
   explicit OpToFuncCallLowering(LLVMTypeConverter &lowering_, StringRef f32Func,
                                 StringRef f64Func)
       : ConvertToLLVMPattern(SourceOp::getOperationName(),
                              lowering_.getDialect()->getContext(), lowering_),
         f32Func(f32Func), f64Func(f64Func) {}
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     using LLVM::LLVMFuncOp;
     using LLVM::LLVMType;
 
     static_assert(
         std::is_base_of<OpTrait::OneResult<SourceOp>, SourceOp>::value,
         "expected single result op");
 
     LLVMType resultType = typeConverter.convertType(op->getResult(0).getType())
                               .template cast<LLVM::LLVMType>();
     LLVMType funcType = getFunctionType(resultType, operands);
     StringRef funcName = getFunctionName(resultType);
     if (funcName.empty())
-      return matchFailure();
+      return failure();
 
     LLVMFuncOp funcOp = appendOrGetFuncOp(funcName, funcType, op);
     auto callOp = rewriter.create<LLVM::CallOp>(
         op->getLoc(), resultType, rewriter.getSymbolRefAttr(funcOp), operands);
     rewriter.replaceOp(op, {callOp.getResult(0)});
-    return matchSuccess();
+    return success();
   }
 
 private:
   LLVM::LLVMType getFunctionType(LLVM::LLVMType resultType,
                                  ArrayRef<Value> operands) const {
     using LLVM::LLVMType;
     SmallVector<LLVMType, 1> operandTypes;
     for (Value operand : operands) {

mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp

struct GPUShuffleOpLowering : public ConvertToLLVMPattern {
   ///     %active_mask = llvm.sub %shl, %one : !llvm.i32
   ///     %mask_and_clamp = llvm.sub %width, %one : !llvm.i32
   ///     %shfl = nvvm.shfl.sync.bfly %active_mask, %value, %offset,
   ///         %mask_and_clamp : !llvm<"{ float, i1 }">
   ///     %shfl_value = llvm.extractvalue %shfl[0 : index] :
   ///         !llvm<"{ float, i1 }">
   ///     %shfl_pred = llvm.extractvalue %shfl[1 : index] :
   ///         !llvm<"{ float, i1 }">
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     Location loc = op->getLoc();
     gpu::ShuffleOpOperandAdaptor adaptor(operands);
 
     auto dialect = typeConverter.getDialect();
     auto valueTy = adaptor.value().getType().cast<LLVM::LLVMType>();
     auto int32Type = LLVM::LLVMType::getInt32Ty(dialect);
Value shfl = rewriter.create<NVVM::ShflBflyOp>(
         loc, resultTy, activeMask, adaptor.value(), adaptor.offset(),
         maskAndClamp, returnValueAndIsValidAttr);
     Value shflValue = rewriter.create<LLVM::ExtractValueOp>(
         loc, valueTy, shfl, rewriter.getIndexArrayAttr(0));
     Value isActiveSrcLane = rewriter.create<LLVM::ExtractValueOp>(
         loc, predTy, shfl, rewriter.getIndexArrayAttr(1));
 
     rewriter.replaceOp(op, {shflValue, isActiveSrcLane});
-    return matchSuccess();
+    return success();
   }
 };
 
 struct GPUFuncOpLowering : ConvertToLLVMPattern {
   explicit GPUFuncOpLowering(LLVMTypeConverter &typeConverter)
       : ConvertToLLVMPattern(gpu::GPUFuncOp::getOperationName(),
                              typeConverter.getDialect()->getContext(),
                              typeConverter) {}
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     assert(operands.empty() && "func op is not expected to have operands");
     auto gpuFuncOp = cast<gpu::GPUFuncOp>(op);
     Location loc = gpuFuncOp.getLoc();
 
     SmallVector<LLVM::GlobalOp, 3> workgroupBuffers;
     workgroupBuffers.reserve(gpuFuncOp.getNumWorkgroupAttributions());
matchAndRewrite(Operation *op, ArrayRef<Value> operands,
 
     // Move the region to the new function, update the entry block signature.
     rewriter.inlineRegionBefore(gpuFuncOp.getBody(), llvmFuncOp.getBody(),
                                 llvmFuncOp.end());
     rewriter.applySignatureConversion(&llvmFuncOp.getBody(),
                                       signatureConversion);
 
     rewriter.eraseOp(gpuFuncOp);
-    return matchSuccess();
+    return success();
   }
 };
 
 struct GPUReturnOpLowering : public ConvertToLLVMPattern {
   GPUReturnOpLowering(LLVMTypeConverter &typeConverter)
       : ConvertToLLVMPattern(gpu::ReturnOp::getOperationName(),
                              typeConverter.getDialect()->getContext(),
                              typeConverter) {}
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     rewriter.replaceOpWithNewOp<LLVM::ReturnOp>(op, operands);
-    return matchSuccess();
+    return success();
   }
 };
 
 /// Import the GPU Ops to NVVM Patterns.
 #include "GPUToNVVM.cpp.inc"
 
 /// A pass that replaces all occurrences of GPU device operations with their
 /// corresponding NVVM equivalent.

mlir/lib/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.cpp

 
 namespace {
 
 /// Pattern to convert a loop::ForOp within kernel functions into spirv::LoopOp.
 class ForOpConversion final : public SPIRVOpLowering<loop::ForOp> {
 public:
   using SPIRVOpLowering<loop::ForOp>::SPIRVOpLowering;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(loop::ForOp forOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
 /// Pattern to convert a loop::IfOp within kernel functions into
 /// spirv::SelectionOp.
 class IfOpConversion final : public SPIRVOpLowering<loop::IfOp> {
 public:
   using SPIRVOpLowering<loop::IfOp>::SPIRVOpLowering;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(loop::IfOp IfOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
 /// Pattern to erase a loop::YieldOp.
 class TerminatorOpConversion final : public SPIRVOpLowering<loop::YieldOp> {
 public:
   using SPIRVOpLowering<loop::YieldOp>::SPIRVOpLowering;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(loop::YieldOp terminatorOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     rewriter.eraseOp(terminatorOp);
-    return matchSuccess();
+    return success();
   }
 };
 
 /// Pattern lowering GPU block/thread size/id to loading SPIR-V invocation
 /// builin variables.
 template <typename SourceOp, spirv::BuiltIn builtin>
 class LaunchConfigConversion : public SPIRVOpLowering<SourceOp> {
 public:
   using SPIRVOpLowering<SourceOp>::SPIRVOpLowering;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(SourceOp op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
 /// This is separate because in Vulkan workgroup size is exposed to shaders via
 /// a constant with WorkgroupSize decoration. So here we cannot generate a
 /// builtin variable; instead the infromation in the `spv.entry_point_abi`
 /// attribute on the surrounding FuncOp is used to replace the gpu::BlockDimOp.
 class WorkGroupSizeConversion : public SPIRVOpLowering<gpu::BlockDimOp> {
 public:
   using SPIRVOpLowering<gpu::BlockDimOp>::SPIRVOpLowering;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(gpu::BlockDimOp op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
 /// Pattern to convert a kernel function in GPU dialect within a spv.module.
 class GPUFuncOpConversion final : public SPIRVOpLowering<gpu::GPUFuncOp> {
 public:
   using SPIRVOpLowering<gpu::GPUFuncOp>::SPIRVOpLowering;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(gpu::GPUFuncOp funcOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override;
 
 private:
   SmallVector<int32_t, 3> workGroupSizeAsInt32;
 };
 
 /// Pattern to convert a gpu.module to a spv.module.
 class GPUModuleConversion final : public SPIRVOpLowering<gpu::GPUModuleOp> {
 public:
   using SPIRVOpLowering<gpu::GPUModuleOp>::SPIRVOpLowering;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(gpu::GPUModuleOp moduleOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
 /// Pattern to convert a gpu.return into a SPIR-V return.
 // TODO: This can go to DRR when GPU return has operands.
 class GPUReturnOpConversion final : public SPIRVOpLowering<gpu::ReturnOp> {
 public:
   using SPIRVOpLowering<gpu::ReturnOp>::SPIRVOpLowering;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(gpu::ReturnOp returnOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
 } // namespace
 
 //===----------------------------------------------------------------------===//
 // loop::ForOp.
 //===----------------------------------------------------------------------===//
 
-PatternMatchResult
+LogicalResult
 ForOpConversion::matchAndRewrite(loop::ForOp forOp, ArrayRef<Value> operands,
                                  ConversionPatternRewriter &rewriter) const {
   // loop::ForOp can be lowered to the structured control flow represented by
   // spirv::LoopOp by making the continue block of the spirv::LoopOp the loop
   // latch and the merge block the exit block. The resulting spirv::LoopOp has a
   // single back edge from the continue to header block, and a single exit from
   // header to merge.
   loop::ForOpOperandAdaptor forOperands(operands);
ForOpConversion::matchAndRewrite(loop::ForOp forOp, ArrayRef<Value> operands,
   rewriter.setInsertionPointToEnd(continueBlock);
 
   // Add the step to the induction variable and branch to the header.
   Value updatedIndVar = rewriter.create<spirv::IAddOp>(
       loc, newIndVar.getType(), newIndVar, forOperands.step());
   rewriter.create<spirv::BranchOp>(loc, header, updatedIndVar);
 
   rewriter.eraseOp(forOp);
-  return matchSuccess();
+  return success();
 }
 
 //===----------------------------------------------------------------------===//
 // loop::IfOp.
 //===----------------------------------------------------------------------===//
 
-PatternMatchResult
+LogicalResult
 IfOpConversion::matchAndRewrite(loop::IfOp ifOp, ArrayRef<Value> operands,
                                 ConversionPatternRewriter &rewriter) const {
   // When lowering `loop::IfOp` we explicitly create a selection header block
   // before the control flow diverges and a merge block where control flow
   // subsequently converges.
   loop::IfOpOperandAdaptor ifOperands(operands);
   auto loc = ifOp.getLoc();
 
IfOpConversion::matchAndRewrite(loop::IfOp ifOp, ArrayRef<Value> operands,
 
   // Create a `spv.BranchConditional` operation for selection header block.
   rewriter.setInsertionPointToEnd(selectionHeaderBlock);
   rewriter.create<spirv::BranchConditionalOp>(loc, ifOperands.condition(),
                                               thenBlock, ArrayRef<Value>(),
                                               elseBlock, ArrayRef<Value>());
 
   rewriter.eraseOp(ifOp);
-  return matchSuccess();
+  return success();
 }
 
 //===----------------------------------------------------------------------===//
 // Builtins.
 //===----------------------------------------------------------------------===//
 
 static Optional<int32_t> getLaunchConfigIndex(Operation *op) {
   auto dimAttr = op->getAttrOfType<StringAttr>("dimension");
   if (!dimAttr) {
     return {};
   }
   if (dimAttr.getValue() == "x") {
     return 0;
   } else if (dimAttr.getValue() == "y") {
     return 1;
   } else if (dimAttr.getValue() == "z") {
     return 2;
   }
   return {};
 }
 
 template <typename SourceOp, spirv::BuiltIn builtin>
-PatternMatchResult LaunchConfigConversion<SourceOp, builtin>::matchAndRewrite(
+LogicalResult LaunchConfigConversion<SourceOp, builtin>::matchAndRewrite(
     SourceOp op, ArrayRef<Value> operands,
     ConversionPatternRewriter &rewriter) const {
   auto index = getLaunchConfigIndex(op);
   if (!index)
-    return this->matchFailure();
+    return failure();
 
   // SPIR-V invocation builtin variables are a vector of type <3xi32>
   auto spirvBuiltin = spirv::getBuiltinVariableValue(op, builtin, rewriter);
   rewriter.replaceOpWithNewOp<spirv::CompositeExtractOp>(
       op, rewriter.getIntegerType(32), spirvBuiltin,
       rewriter.getI32ArrayAttr({index.getValue()}));
-  return this->matchSuccess();
+  return success();
 }
 
-PatternMatchResult WorkGroupSizeConversion::matchAndRewrite(
+LogicalResult WorkGroupSizeConversion::matchAndRewrite(
     gpu::BlockDimOp op, ArrayRef<Value> operands,
     ConversionPatternRewriter &rewriter) const {
   auto index = getLaunchConfigIndex(op);
   if (!index)
-    return matchFailure();
+    return failure();
 
   auto workGroupSizeAttr = spirv::lookupLocalWorkGroupSize(op);
   auto val = workGroupSizeAttr.getValue<int32_t>(index.getValue());
   auto convertedType = typeConverter.convertType(op.getResult().getType());
   if (!convertedType)
-    return matchFailure();
+    return failure();
   rewriter.replaceOpWithNewOp<spirv::ConstantOp>(
       op, convertedType, IntegerAttr::get(convertedType, val));
-  return matchSuccess();
+  return success();
 }
 
 //===----------------------------------------------------------------------===//
 // GPUFuncOp
 //===----------------------------------------------------------------------===//
 
 // Legalizes a GPU function as an entry SPIR-V function.
 static spirv::FuncOp
rewriter.inlineRegionBefore(funcOp.getBody(), newFuncOp.getBody(),
                               newFuncOp.end());
   rewriter.applySignatureConversion(&newFuncOp.getBody(), signatureConverter);
   rewriter.eraseOp(funcOp);
 
   spirv::setABIAttrs(newFuncOp, entryPointInfo, argABIInfo);
   return newFuncOp;
 }
 
-PatternMatchResult GPUFuncOpConversion::matchAndRewrite(
+LogicalResult GPUFuncOpConversion::matchAndRewrite(
     gpu::GPUFuncOp funcOp, ArrayRef<Value> operands,
     ConversionPatternRewriter &rewriter) const {
   if (!gpu::GPUDialect::isKernel(funcOp))
-    return matchFailure();
+    return failure();
 
   SmallVector<spirv::InterfaceVarABIAttr, 4> argABI;
   for (auto argNum : llvm::seq<unsigned>(0, funcOp.getNumArguments())) {
     argABI.push_back(spirv::getInterfaceVarABIAttr(
         0, argNum, spirv::StorageClass::StorageBuffer, rewriter.getContext()));
   }
 
   auto entryPointAttr = spirv::lookupEntryPointABI(funcOp);
   if (!entryPointAttr) {
     funcOp.emitRemark("match failure: missing 'spv.entry_point_abi' attribute");
-    return matchFailure();
+    return failure();
   }
   spirv::FuncOp newFuncOp = lowerAsEntryFunction(
       funcOp, typeConverter, rewriter, entryPointAttr, argABI);
   if (!newFuncOp)
-    return matchFailure();
+    return failure();
   newFuncOp.removeAttr(Identifier::get(gpu::GPUDialect::getKernelFuncAttrName(),
                                        rewriter.getContext()));
-  return matchSuccess();
+  return success();
 }
 
 //===----------------------------------------------------------------------===//
 // ModuleOp with gpu.module.
 //===----------------------------------------------------------------------===//
 
-PatternMatchResult GPUModuleConversion::matchAndRewrite(
+LogicalResult GPUModuleConversion::matchAndRewrite(
     gpu::GPUModuleOp moduleOp, ArrayRef<Value> operands,
     ConversionPatternRewriter &rewriter) const {
   auto spvModule = rewriter.create<spirv::ModuleOp>(
       moduleOp.getLoc(), spirv::AddressingModel::Logical,
       spirv::MemoryModel::GLSL450);
 
   // Move the region from the module op into the SPIR-V module.
   Region &spvModuleRegion = spvModule.body();
   rewriter.inlineRegionBefore(moduleOp.body(), spvModuleRegion,
                               spvModuleRegion.begin());
   // The spv.module build method adds a block with a terminator. Remove that
   // block. The terminator of the module op in the remaining block will be
   // legalized later.
   spvModuleRegion.back().erase();
   rewriter.eraseOp(moduleOp);
-  return matchSuccess();
+  return success();
 }
 
 //===----------------------------------------------------------------------===//
 // GPU return inside kernel functions to SPIR-V return.
 //===----------------------------------------------------------------------===//
 
-PatternMatchResult GPUReturnOpConversion::matchAndRewrite(
+LogicalResult GPUReturnOpConversion::matchAndRewrite(
     gpu::ReturnOp returnOp, ArrayRef<Value> operands,
     ConversionPatternRewriter &rewriter) const {
   if (!operands.empty())
-    return matchFailure();
+    return failure();
 
   rewriter.replaceOpWithNewOp<spirv::ReturnOp>(returnOp);
-  return matchSuccess();
+  return success();
 }
 
 //===----------------------------------------------------------------------===//
 // GPU To SPIRV Patterns.
 //===----------------------------------------------------------------------===//
 
 namespace {
 #include "GPUToSPIRV.cpp.inc"

mlir/lib/Conversion/LinalgToLLVM/LinalgToLLVM.cpp

 };
 
 // RangeOp creates a new range descriptor.
 class RangeOpConversion : public ConvertToLLVMPattern {
 public:
   explicit RangeOpConversion(MLIRContext *context, LLVMTypeConverter &lowering_)
       : ConvertToLLVMPattern(RangeOp::getOperationName(), context, lowering_) {}
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto rangeOp = cast<RangeOp>(op);
     auto rangeDescriptorTy =
         convertRangeType(rangeOp.getType().cast<RangeType>(), typeConverter);
 
     edsc::ScopedContext context(rewriter, op->getLoc());
 
     // Fill in an aggregate value of the descriptor.
     RangeOpOperandAdaptor adaptor(operands);
     Value desc = llvm_undef(rangeDescriptorTy);
     desc = llvm_insertvalue(desc, adaptor.min(), rewriter.getI64ArrayAttr(0));
     desc = llvm_insertvalue(desc, adaptor.max(), rewriter.getI64ArrayAttr(1));
     desc = llvm_insertvalue(desc, adaptor.step(), rewriter.getI64ArrayAttr(2));
     rewriter.replaceOp(op, desc);
-    return matchSuccess();
+    return success();
   }
 };
 
 // ReshapeOp creates a new view descriptor of the proper rank.
 // For now, the only conversion supported is for target MemRef with static sizes
 // and strides.
 class ReshapeOpConversion : public ConvertToLLVMPattern {
 public:
   explicit ReshapeOpConversion(MLIRContext *context,
                                LLVMTypeConverter &lowering_)
       : ConvertToLLVMPattern(ReshapeOp::getOperationName(), context,
                              lowering_) {}
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto reshapeOp = cast<ReshapeOp>(op);
     MemRefType dstType = reshapeOp.getResult().getType().cast<MemRefType>();
 
     if (!dstType.hasStaticShape())
-      return matchFailure();
+      return failure();
 
     int64_t offset;
     SmallVector<int64_t, 4> strides;
     auto res = getStridesAndOffset(dstType, strides, offset);
     if (failed(res) || llvm::any_of(strides, [](int64_t val) {
           return ShapedType::isDynamicStrideOrOffset(val);
         }))
-      return matchFailure();
+      return failure();
 
     edsc::ScopedContext context(rewriter, op->getLoc());
     ReshapeOpOperandAdaptor adaptor(operands);
     BaseViewConversionHelper baseDesc(adaptor.view());
     BaseViewConversionHelper desc(typeConverter.convertType(dstType));
     desc.setAllocatedPtr(baseDesc.allocatedPtr());
     desc.setAlignedPtr(baseDesc.alignedPtr());
     desc.setOffset(baseDesc.offset());
     for (auto en : llvm::enumerate(dstType.getShape()))
       desc.setConstantSize(en.index(), en.value());
     for (auto en : llvm::enumerate(strides))
       desc.setConstantStride(en.index(), en.value());
     rewriter.replaceOp(op, {desc});
-    return matchSuccess();
+    return success();
   }
 };
 
 /// Conversion pattern that transforms a linalg.slice op into:
 ///   1. An "undef" value for the ViewDescriptor.
 ///   2. Updates to the ViewDescriptor to introduce the data ptr, offset, size
 ///      and stride corresponding to the region of memory within the bounds of
 ///      the parent view.
 /// The linalg.slice op is replaced by the alloca'ed pointer.
 class SliceOpConversion : public ConvertToLLVMPattern {
 public:
   explicit SliceOpConversion(MLIRContext *context, LLVMTypeConverter &lowering_)
       : ConvertToLLVMPattern(SliceOp::getOperationName(), context, lowering_) {}
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     edsc::ScopedContext context(rewriter, op->getLoc());
     SliceOpOperandAdaptor adaptor(operands);
     BaseViewConversionHelper baseDesc(adaptor.view());
 
     auto sliceOp = cast<SliceOp>(op);
     auto memRefType = sliceOp.getBaseViewType();
matchAndRewrite(Operation *op, ArrayRef<Value> operands,
     desc.setAllocatedPtr(baseDesc.allocatedPtr());
     desc.setAlignedPtr(baseDesc.alignedPtr());
 
     // Insert base offset.
     desc.setOffset(baseOffset);
 
     // Corner case, no sizes or strides: early return the descriptor.
     if (sliceOp.getShapedType().getRank() == 0)
-      return rewriter.replaceOp(op, {desc}), matchSuccess();
+      return rewriter.replaceOp(op, {desc}), success();
 
     Value zero = llvm_constant(
         int64Ty, rewriter.getIntegerAttr(rewriter.getIndexType(), 0));
     // Compute and insert view sizes (max - min along the range) and strides.
     // Skip the non-range operands as they will be projected away from the view.
     int numNewDims = 0;
     for (auto en : llvm::enumerate(sliceOp.indexings())) {
       Value indexing = en.value();
for (auto en : llvm::enumerate(sliceOp.indexings())) {
         Value stride = llvm_mul(strides[rank], step);
         desc.setSize(numNewDims, size);
         desc.setStride(numNewDims, stride);
         ++numNewDims;
       }
     }
 
     rewriter.replaceOp(op, {desc});
-    return matchSuccess();
+    return success();
   }
 };
 
 /// Conversion pattern that transforms a linalg.transpose op into:
 ///   1. A function entry `alloca` operation to allocate a ViewDescriptor.
 ///   2. A load of the ViewDescriptor from the pointer allocated in 1.
 ///   3. Updates to the ViewDescriptor to introduce the data ptr, offset, size
 ///      and stride. Size and stride are permutations of the original values.
 ///   4. A store of the resulting ViewDescriptor to the alloca'ed pointer.
 /// The linalg.transpose op is replaced by the alloca'ed pointer.
 class TransposeOpConversion : public ConvertToLLVMPattern {
 public:
   explicit TransposeOpConversion(MLIRContext *context,
                                  LLVMTypeConverter &lowering_)
       : ConvertToLLVMPattern(TransposeOp::getOperationName(), context,
                              lowering_) {}
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     // Initialize the common boilerplate and alloca at the top of the FuncOp.
     edsc::ScopedContext context(rewriter, op->getLoc());
     TransposeOpOperandAdaptor adaptor(operands);
     BaseViewConversionHelper baseDesc(adaptor.view());
 
     auto transposeOp = cast<TransposeOp>(op);
     // No permutation, early exit.
     if (transposeOp.permutation().isIdentity())
-      return rewriter.replaceOp(op, {baseDesc}), matchSuccess();
+      return rewriter.replaceOp(op, {baseDesc}), success();
 
     BaseViewConversionHelper desc(
         typeConverter.convertType(transposeOp.getShapedType()));
 
     // Copy the base and aligned pointers from the old descriptor to the new
     // one.
     desc.setAllocatedPtr(baseDesc.allocatedPtr());
     desc.setAlignedPtr(baseDesc.alignedPtr());
 
     // Copy the offset pointer from the old descriptor to the new one.
     desc.setOffset(baseDesc.offset());
 
     // Iterate over the dimensions and apply size/stride permutation.
     for (auto en : llvm::enumerate(transposeOp.permutation().getResults())) {
       int sourcePos = en.index();
       int targetPos = en.value().cast<AffineDimExpr>().getPosition();
       desc.setSize(targetPos, baseDesc.size(sourcePos));
       desc.setStride(targetPos, baseDesc.stride(sourcePos));
     }
 
     rewriter.replaceOp(op, {desc});
-    return matchSuccess();
+    return success();
   }
 };
 
 // YieldOp produces and LLVM::ReturnOp.
 class YieldOpConversion : public ConvertToLLVMPattern {
 public:
   explicit YieldOpConversion(MLIRContext *context, LLVMTypeConverter &lowering_)
       : ConvertToLLVMPattern(YieldOp::getOperationName(), context, lowering_) {}
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     rewriter.replaceOpWithNewOp<LLVM::ReturnOp>(op, operands);
-    return matchSuccess();
+    return success();
   }
 };
 } // namespace
 
 template <typename LinalgOp>
 static SmallVector<Type, 4> ExtractOperandTypes(Operation *op) {
   return SmallVector<Type, 4>{op->getOperandTypes()};
 }

 // `LinalgOp::getLibraryCallName()` function.
 // The implementation of the function can be either in the same module or in an
 // externally linked library.
 template <typename LinalgOp>
 class LinalgOpConversion : public OpRewritePattern<LinalgOp> {
 public:
   using OpRewritePattern<LinalgOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(LinalgOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(LinalgOp op,
     auto libraryCallName = getLibraryCallSymbolRef<LinalgOp>(op, rewriter);
     if (!libraryCallName)
-      return this->matchFailure();
+      return failure();
 
     rewriter.replaceOpWithNewOp<mlir::CallOp>(
         op, libraryCallName.getValue(), ArrayRef<Type>{}, op.getOperands());
-    return this->matchSuccess();
+    return success();
   }
 };
 
 /// Conversion pattern specialization for CopyOp. This kicks in when both input
 /// and output permutations are left unspecified or are the identity.
 template <> class LinalgOpConversion<CopyOp> : public OpRewritePattern<CopyOp> {
 public:
   using OpRewritePattern<CopyOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(CopyOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(CopyOp op,
     auto inputPerm = op.inputPermutation();
     if (inputPerm.hasValue() && !inputPerm->isIdentity())
-      return matchFailure();
+      return failure();
     auto outputPerm = op.outputPermutation();
     if (outputPerm.hasValue() && !outputPerm->isIdentity())
-      return matchFailure();
+      return failure();
 
     auto libraryCallName = getLibraryCallSymbolRef<CopyOp>(op, rewriter);
     if (!libraryCallName)
-      return matchFailure();
+      return failure();
 
     rewriter.replaceOpWithNewOp<mlir::CallOp>(
         op, libraryCallName.getValue(), ArrayRef<Type>{}, op.getOperands());
-    return matchSuccess();
+    return success();
   }
 };
 
 /// Conversion pattern specialization for IndexedGenericOp.
 template <>
 class LinalgOpConversion<IndexedGenericOp>
     : public OpRewritePattern<IndexedGenericOp> {
 public:
   using OpRewritePattern<IndexedGenericOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(IndexedGenericOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(IndexedGenericOp op,
     auto libraryCallName =
         getLibraryCallSymbolRef<IndexedGenericOp>(op, rewriter);
     if (!libraryCallName)
-      return this->matchFailure();
+      return failure();
 
     // TODO(pifon, ntv): Use induction variables values instead of zeros, when
     // IndexedGenericOp is tiled.
     auto zero = rewriter.create<mlir::ConstantOp>(
         op.getLoc(), rewriter.getIntegerAttr(rewriter.getIndexType(), 0));
     auto indexedGenericOp = cast<IndexedGenericOp>(op);
     auto numLoops = indexedGenericOp.getNumLoops();
     SmallVector<Value, 4> operands;
     operands.reserve(numLoops + op.getNumOperands());
     for (unsigned i = 0; i < numLoops; ++i) {
       operands.push_back(zero);
     }
     for (auto operand : op.getOperands()) {
       operands.push_back(operand);
     }
     rewriter.replaceOpWithNewOp<mlir::CallOp>(op, libraryCallName.getValue(),
                                               ArrayRef<Type>{}, operands);
-    return this->matchSuccess();
+    return success();
   }
 };
 
 /// A non-conversion rewrite pattern kicks in to convert CopyOp with
 /// permutations into a sequence of TransposeOp and permutation-free CopyOp.
 /// This interplays together with TransposeOpConversion and
 /// LinalgConversion<CopyOp> to create a path to the LLVM dialect.
 class CopyTransposeConversion : public OpRewritePattern<CopyOp> {
 public:
   using OpRewritePattern<CopyOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(CopyOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(CopyOp op,
     Value in = op.input(), out = op.output();
 
     // If either inputPerm or outputPerm are non-identities, insert transposes.
     auto inputPerm = op.inputPermutation();
     if (inputPerm.hasValue() && !inputPerm->isIdentity())
       in = rewriter.create<linalg::TransposeOp>(op.getLoc(), in,
                                                 AffineMapAttr::get(*inputPerm));
     auto outputPerm = op.outputPermutation();
     if (outputPerm.hasValue() && !outputPerm->isIdentity())
       out = rewriter.create<linalg::TransposeOp>(
           op.getLoc(), out, AffineMapAttr::get(*outputPerm));
 
     // If nothing was transposed, fail and let the conversion kick in.
     if (in == op.input() && out == op.output())
-      return matchFailure();
+      return failure();
 
     rewriter.replaceOpWithNewOp<CopyOp>(op, in, out);
-    return matchSuccess();
+    return success();
   }
 };
 
 /// Populate the given list with patterns that convert from Linalg to Standard.
 static void
 populateLinalgToStandardConversionPatterns(OwningRewritePatternList &patterns,
                                            MLIRContext *ctx) {
   // TODO(ntv) ConvOp conversion needs to export a descriptor with relevant

mlir/lib/Conversion/LinalgToSPIRV/LinalgToSPIRV.cpp

 public:
   using SPIRVOpLowering<linalg::GenericOp>::SPIRVOpLowering;
 
   /// Matches the given linalg.generic op as performing reduction and returns
   /// the binary op kind if successful.
   static Optional<linalg::RegionMatcher::BinaryOpKind>
   matchAsPerformingReduction(linalg::GenericOp genericOp);
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(linalg::GenericOp genericOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
 } // namespace
 
 Optional<linalg::RegionMatcher::BinaryOpKind>
 SingleWorkgroupReduction::matchAsPerformingReduction(
SingleWorkgroupReduction::matchAsPerformingReduction(
   // The indexing map for the input should be `(i) -> (0)`.
   if (outputMap.getValue() !=
       AffineMap::get(1, 0, {getAffineConstantExpr(0, op->getContext())}))
     return llvm::None;
 
   return linalg::RegionMatcher::matchAsScalarBinaryOp(genericOp);
 }
 
-PatternMatchResult SingleWorkgroupReduction::matchAndRewrite(
+LogicalResult SingleWorkgroupReduction::matchAndRewrite(
     linalg::GenericOp genericOp, ArrayRef<Value> operands,
     ConversionPatternRewriter &rewriter) const {
   Operation *op = genericOp.getOperation();
   auto originalInputType = op->getOperand(0).getType().cast<MemRefType>();
   auto originalOutputType = op->getOperand(1).getType().cast<MemRefType>();
 
   auto binaryOpKind = matchAsPerformingReduction(genericOp);
   if (!binaryOpKind)
-    return matchFailure();
+    return failure();
 
   // Query the shader interface for local workgroup size to make sure the
   // invocation configuration fits with the input memref's shape.
   DenseIntElementsAttr localSize = spirv::lookupLocalWorkGroupSize(genericOp);
   if (!localSize)
-    return matchFailure();
+    return failure();
 
   if ((*localSize.begin()).getSExtValue() != originalInputType.getDimSize(0))
-    return matchFailure();
+    return failure();
   if (llvm::any_of(llvm::drop_begin(localSize.getIntValues(), 1),
                    [](const APInt &size) { return !size.isOneValue(); }))
-    return matchFailure();
+    return failure();
 
   // TODO(antiagainst): Query the target environment to make sure the current
   // workload fits in a local workgroup.
 
   Value convertedInput = operands[0], convertedOutput = operands[1];
   Location loc = genericOp.getLoc();
 
   // Get the invocation ID.
#define CREATE_ATOMIC_BIN_OP(opKind, spvOp)                                    \
   } break
     switch (*binaryOpKind) { CREATE_ATOMIC_BIN_OP(IAdd, AtomicIAddOp); }
 #undef CREATE_ATOMIC_BIN_OP
   };
 
   spirv::SelectionOp::createIfThen(loc, condition, createAtomicOp, &rewriter);
 
   rewriter.eraseOp(genericOp);
-  return matchSuccess();
+  return success();
 }
 
 //===----------------------------------------------------------------------===//
 // Pattern population
 //===----------------------------------------------------------------------===//
 
 void mlir::populateLinalgToSPIRVPatterns(MLIRContext *context,
                                          SPIRVTypeConverter &typeConverter,
                                          OwningRewritePatternList &patterns) {
   patterns.insert<SingleWorkgroupReduction>(context, typeConverter);
 }

mlir/lib/Conversion/LoopToStandard/ConvertLoopToStandard.cpp

 //      | end:                           |
 //      |   <code after the ForOp>       |
 //      |   <%init visible by dominance> |
 //      +--------------------------------+
 //
 struct ForLowering : public OpRewritePattern<ForOp> {
   using OpRewritePattern<ForOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(ForOp forOp,
-                                     PatternRewriter &rewriter) const override;
+  LogicalResult matchAndRewrite(ForOp forOp,
 };
 
 // Create a CFG subgraph for the loop.if operation (including its "then" and
 // optional "else" operation blocks).  We maintain the invariants that the
 // subgraph has a single entry and a single exit point, and that the entry/exit
 // blocks are respectively the first/last block of the enclosing region. The
 // operations following the loop.if are split into a continuation (subgraph
 // exit) block. The condition is lowered to a chain of blocks that implement the

 //      +--------------------------------+
 //      | continue:                      |
 //      |   <code after the IfOp>        |
 //      +--------------------------------+
 //
 struct IfLowering : public OpRewritePattern<IfOp> {
   using OpRewritePattern<IfOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(IfOp ifOp,
-                                     PatternRewriter &rewriter) const override;
+  LogicalResult matchAndRewrite(IfOp ifOp,
 };
 
 struct ParallelLowering : public OpRewritePattern<mlir::loop::ParallelOp> {
   using OpRewritePattern<mlir::loop::ParallelOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(mlir::loop::ParallelOp parallelOp,
-                                     PatternRewriter &rewriter) const override;
+  LogicalResult matchAndRewrite(mlir::loop::ParallelOp parallelOp,
 };
 } // namespace
 
-PatternMatchResult
-ForLowering::matchAndRewrite(ForOp forOp, PatternRewriter &rewriter) const {
+LogicalResult ForLowering::matchAndRewrite(ForOp forOp,
+                                           PatternRewriter &rewriter) const {
   Location loc = forOp.getLoc();
 
   // Start by splitting the block containing the 'loop.for' into two parts.
   // The part before will get the init code, the part after will be the end
   // point.
   auto *initBlock = rewriter.getInsertionBlock();
   auto initPosition = rewriter.getInsertionPoint();
   auto *endBlock = rewriter.splitBlock(initBlock, initPosition);
LogicalResult ForLowering::matchAndRewrite(ForOp forOp,
   // Append the induction variable stepping logic to the last body block and
   // branch back to the condition block. Loop-carried values are taken from
   // operands of the loop terminator.
   Operation *terminator = lastBodyBlock->getTerminator();
   rewriter.setInsertionPointToEnd(lastBodyBlock);
   auto step = forOp.step();
   auto stepped = rewriter.create<AddIOp>(loc, iv, step).getResult();
   if (!stepped)
-    return matchFailure();
+    return failure();
 
   SmallVector<Value, 8> loopCarried;
   loopCarried.push_back(stepped);
   loopCarried.append(terminator->operand_begin(), terminator->operand_end());
   rewriter.create<BranchOp>(loc, conditionBlock, loopCarried);
   rewriter.eraseOp(terminator);
 
   // Compute loop bounds before branching to the condition.
   rewriter.setInsertionPointToEnd(initBlock);
   Value lowerBound = forOp.lowerBound();
   Value upperBound = forOp.upperBound();
   if (!lowerBound || !upperBound)
-    return matchFailure();
+    return failure();
 
   // The initial values of loop-carried values is obtained from the operands
   // of the loop operation.
   SmallVector<Value, 8> destOperands;
   destOperands.push_back(lowerBound);
   auto iterOperands = forOp.getIterOperands();
   destOperands.append(iterOperands.begin(), iterOperands.end());
   rewriter.create<BranchOp>(loc, conditionBlock, destOperands);
 
   // With the body block done, we can fill in the condition block.
   rewriter.setInsertionPointToEnd(conditionBlock);
   auto comparison =
       rewriter.create<CmpIOp>(loc, CmpIPredicate::slt, iv, upperBound);
 
   rewriter.create<CondBranchOp>(loc, comparison, firstBodyBlock,
                                 ArrayRef<Value>(), endBlock, ArrayRef<Value>());
   // The result of the loop operation is the values of the condition block
   // arguments except the induction variable on the last iteration.
   rewriter.replaceOp(forOp, conditionBlock->getArguments().drop_front());
-  return matchSuccess();
+  return success();
 }
 
-PatternMatchResult
-IfLowering::matchAndRewrite(IfOp ifOp, PatternRewriter &rewriter) const {
+LogicalResult IfLowering::matchAndRewrite(IfOp ifOp,
+                                          PatternRewriter &rewriter) const {
   auto loc = ifOp.getLoc();
 
   // Start by splitting the block containing the 'loop.if' into two parts.
   // The part before will contain the condition, the part after will be the
   // continuation point.
   auto *condBlock = rewriter.getInsertionBlock();
   auto opPosition = rewriter.getInsertionPoint();
   auto *continueBlock = rewriter.splitBlock(condBlock, opPosition);
LogicalResult IfLowering::matchAndRewrite(IfOp ifOp,
 
   rewriter.setInsertionPointToEnd(condBlock);
   rewriter.create<CondBranchOp>(loc, ifOp.condition(), thenBlock,
                                 /*trueArgs=*/ArrayRef<Value>(), elseBlock,
                                 /*falseArgs=*/ArrayRef<Value>());
 
   // Ok, we're done!
   rewriter.eraseOp(ifOp);
-  return matchSuccess();
+  return success();
 }
 
-PatternMatchResult
+LogicalResult
 ParallelLowering::matchAndRewrite(ParallelOp parallelOp,
                                   PatternRewriter &rewriter) const {
   Location loc = parallelOp.getLoc();
   BlockAndValueMapping mapping;
 
   // For a parallel loop, we essentially need to create an n-dimensional loop
   // nest. We do this by translating to loop.for ops and have those lowered in
   // a further rewrite. If a parallel loop contains reductions (and thus returns
ParallelLowering::matchAndRewrite(ParallelOp parallelOp,
   }
 
   rewriter.setInsertionPointToEnd(rewriter.getInsertionBlock());
   rewriter.replaceOpWithNewOp<YieldOp>(
       rewriter.getInsertionBlock()->getTerminator(), yieldOperands);
 
   rewriter.replaceOp(parallelOp, loopResults);
 
-  return matchSuccess();
+  return success();
 }
 
 void mlir::populateLoopToStdConversionPatterns(
     OwningRewritePatternList &patterns, MLIRContext *ctx) {
   patterns.insert<ForLowering, IfLowering, ParallelLowering>(ctx);
 }
 
 void LoopToStandardPass::runOnOperation() {

mlir/lib/Conversion/LoopsToGPU/LoopsToGPU.cpp

LogicalResult mlir::convertLoopToGPULaunch(loop::ForOp forOp,
                                            ArrayRef<Value> workGroupSizes) {
   return ::convertLoopToGPULaunch(forOp, numWorkGroups, workGroupSizes);
 }
 
 namespace {
 struct ParallelToGpuLaunchLowering : public OpRewritePattern<ParallelOp> {
   using OpRewritePattern<ParallelOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(ParallelOp parallelOp,
-                                     PatternRewriter &rewriter) const override;
+  LogicalResult matchAndRewrite(ParallelOp parallelOp,
 };
 
 struct MappingAnnotation {
   unsigned processor;
   AffineMap indexMap;
   AffineMap boundMap;
 };
 

 /// launch dimension for the hardware id from the number of iterations the
 /// mapped loop is performing. Note that the number of iterations might be
 /// imprecise if the corresponding loop-bounds are loop-dependent. In such case,
 /// the hardware id might iterate over additional indices. The transformation
 /// caters for this by predicating the created sequence of instructions on
 /// the actual loop bound. This only works if an static upper bound for the
 /// dynamic loop bound can be defived, currently via analyzing `affine.min`
 /// operations.
-PatternMatchResult
+LogicalResult
 ParallelToGpuLaunchLowering::matchAndRewrite(ParallelOp parallelOp,
                                              PatternRewriter &rewriter) const {
   // Create a launch operation. We start with bound one for all grid/block
   // sizes. Those will be refined later as we discover them from mappings.
   Location loc = parallelOp.getLoc();
   Value constantOne = rewriter.create<ConstantIndexOp>(parallelOp.getLoc(), 1);
   gpu::LaunchOp launchOp = rewriter.create<gpu::LaunchOp>(
       parallelOp.getLoc(), constantOne, constantOne, constantOne, constantOne,
       constantOne, constantOne);
   rewriter.setInsertionPointToEnd(&launchOp.body().front());
   rewriter.create<gpu::TerminatorOp>(loc);
   rewriter.setInsertionPointToStart(&launchOp.body().front());
 
   BlockAndValueMapping cloningMap;
   llvm::DenseMap<int, Value> launchBounds;
   SmallVector<Operation *, 16> worklist;
   if (failed(processParallelLoop(parallelOp, launchOp, cloningMap, worklist,
                                  launchBounds, rewriter)))
-    return matchFailure();
+    return failure();
 
   // Whether we have seen any side-effects. Reset when leaving an inner scope.
   bool seenSideeffects = false;
   // Whether we have left a nesting scope (and hence are no longer innermost).
   bool leftNestingScope = false;
   while (!worklist.empty()) {
     Operation *op = worklist.pop_back_val();
     // Now walk over the body and clone it.
     // TODO: This is only correct if there either is no further loop.parallel
     //       nested or this code is side-effect free. Otherwise we might need
     //       predication. We are overly consertaive for now and only allow
     //       side-effects in the innermost scope.
     if (auto nestedParallel = dyn_cast<ParallelOp>(op)) {
       // Before entering a nested scope, make sure there have been no
       // sideeffects until now.
       if (seenSideeffects)
-        return matchFailure();
+        return failure();
       // A nested loop.parallel needs insertion of code to compute indices.
       // Insert that now. This will also update the worklist with the loops
       // body.
       if (failed(processParallelLoop(nestedParallel, launchOp, cloningMap,
                                      worklist, launchBounds, rewriter)))
-        return matchFailure();
+        return failure();
     } else if (op == launchOp.getOperation()) {
       // Found our sentinel value. We have finished the operations from one
       // nesting level, pop one level back up.
       auto parent = rewriter.getInsertionPoint()->getParentOp();
       rewriter.setInsertionPointAfter(parent);
       leftNestingScope = true;
       seenSideeffects = false;
     } else {
       // Otherwise we copy it over.
       Operation *clone = rewriter.clone(*op, cloningMap);
       cloningMap.map(op->getResults(), clone->getResults());
       // Check for side effects.
       // TODO: Handle region side effects properly.
       seenSideeffects |= !MemoryEffectOpInterface::hasNoEffect(clone) ||
                          clone->getNumRegions() != 0;
       // If we are no longer in the innermost scope, sideeffects are disallowed.
       if (seenSideeffects && leftNestingScope)
-        return matchFailure();
+        return failure();
     }
   }
 
   // Now that we succeeded creating the launch operation, also update the
   // bounds.
   for (auto bound : launchBounds)
     launchOp.setOperand(std::get<0>(bound), std::get<1>(bound));
 
   rewriter.eraseOp(parallelOp);
-  return matchSuccess();
+  return success();
 }
 
 void mlir::populateParallelLoopToGPUPatterns(OwningRewritePatternList &patterns,
                                              MLIRContext *ctx) {
   patterns.insert<ParallelToGpuLaunchLowering>(ctx);
 }

mlir/lib/Conversion/StandardToLLVM/ConvertStandardToLLVM.cpp

 /// FuncOp legalization pattern that converts MemRef arguments to pointers to
 /// MemRef descriptors (LLVM struct data types) containing all the MemRef type
 /// information.
 struct FuncOpConversion : public FuncOpConversionBase {
   FuncOpConversion(LLVM::LLVMDialect &dialect, LLVMTypeConverter &converter,
                    bool emitCWrappers)
       : FuncOpConversionBase(dialect, converter), emitWrappers(emitCWrappers) {}
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto funcOp = cast<FuncOp>(op);
 
     auto newFuncOp = convertFuncOpToLLVMFuncOp(funcOp, rewriter);
     if (emitWrappers) {
       if (newFuncOp.isExternal())
         wrapExternalFunction(rewriter, op->getLoc(), typeConverter, funcOp,
                              newFuncOp);
       else
         wrapForExternalCallers(rewriter, op->getLoc(), typeConverter, funcOp,
                                newFuncOp);
     }
 
     rewriter.eraseOp(op);
-    return matchSuccess();
+    return success();
   }
 
 private:
   /// If true, also create the adaptor functions having signatures compatible
   /// with those produced by clang.
   const bool emitWrappers;
 };
 
 /// FuncOp legalization pattern that converts MemRef arguments to bare pointers
 /// to the MemRef element type. This will impact the calling convention and ABI.
 struct BarePtrFuncOpConversion : public FuncOpConversionBase {
   using FuncOpConversionBase::FuncOpConversionBase;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto funcOp = cast<FuncOp>(op);
 
     // Store the positions and type of memref-typed arguments so that we can
     // promote them to MemRef descriptor structs at the beginning of the
     // function.
     SmallVector<UnsignedTypePair, 4> promotedArgsInfo;
     getMemRefArgIndicesAndTypes(funcOp.getType(), promotedArgsInfo);
 
     auto newFuncOp = convertFuncOpToLLVMFuncOp(funcOp, rewriter);
     if (newFuncOp.getBody().empty()) {
       rewriter.eraseOp(op);
-      return matchSuccess();
+      return success();
     }
 
     // Promote bare pointers from MemRef arguments to a MemRef descriptor struct
     // at the beginning of the function so that all the MemRefs in the function
     // have a uniform representation.
     Block *firstBlock = &newFuncOp.getBody().front();
     rewriter.setInsertionPoint(firstBlock, firstBlock->begin());
     auto funcLoc = funcOp.getLoc();
for (const auto &argInfo : promotedArgsInfo) {
         rewriter.replaceUsesOfBlockArgument(arg, placeHolder);
         auto desc = MemRefDescriptor::fromStaticShape(
             rewriter, funcLoc, typeConverter, memrefType, arg);
         rewriter.replaceOp(placeHolder.getDefiningOp(), {desc});
       }
     }
 
     rewriter.eraseOp(op);
-    return matchSuccess();
+    return success();
   }
 };
 
 //////////////// Support for Lowering operations on n-D vectors ////////////////
 namespace {
 // Helper struct to "unroll" operations on n-D vectors in terms of operations on
 // 1-D LLVM vectors.
 struct NDVectorTypeInfo {

 // LLVM Dialect Ops.
 template <typename SourceOp, typename TargetOp>
 struct OneToOneLLVMOpLowering : public LLVMLegalizationPattern<SourceOp> {
   using LLVMLegalizationPattern<SourceOp>::LLVMLegalizationPattern;
   using Super = OneToOneLLVMOpLowering<SourceOp, TargetOp>;
 
   // Convert the type of the result to an LLVM type, pass operands as is,
   // preserve attributes.
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     unsigned numResults = op->getNumResults();
 
     Type packedType;
     if (numResults != 0) {
       packedType =
           this->typeConverter.packFunctionResults(op->getResultTypes());
       if (!packedType)
-        return this->matchFailure();
+        return failure();
     }
 
     auto newOp = rewriter.create<TargetOp>(op->getLoc(), packedType, operands,
                                            op->getAttrs());
 
     // If the operation produced 0 or 1 result, return them immediately.
     if (numResults == 0)
-      return rewriter.eraseOp(op), this->matchSuccess();
+      return rewriter.eraseOp(op), success();
     if (numResults == 1)
       return rewriter.replaceOp(op, newOp.getOperation()->getResult(0)),
-             this->matchSuccess();
+             success();
 
     // Otherwise, it had been converted to an operation producing a structure.
     // Extract individual results from the structure and return them as list.
     SmallVector<Value, 4> results;
     results.reserve(numResults);
     for (unsigned i = 0; i < numResults; ++i) {
       auto type = this->typeConverter.convertType(op->getResult(i).getType());
       results.push_back(rewriter.create<LLVM::ExtractValueOp>(
           op->getLoc(), type, newOp.getOperation()->getResult(0),
           rewriter.getI64ArrayAttr(i)));
     }
     rewriter.replaceOp(op, results);
-    return this->matchSuccess();
+    return success();
   }
 };
 
 template <typename SourceOp, unsigned OpCount>
 struct OpCountValidator {
   static_assert(
       std::is_base_of<
           typename OpTrait::NOperands<OpCount>::template Impl<SourceOp>,

 // types.
 template <typename SourceOp, typename TargetOp, unsigned OpCount>
 struct NaryOpLLVMOpLowering : public LLVMLegalizationPattern<SourceOp> {
   using LLVMLegalizationPattern<SourceOp>::LLVMLegalizationPattern;
   using Super = NaryOpLLVMOpLowering<SourceOp, TargetOp, OpCount>;
 
   // Convert the type of the result to an LLVM type, pass operands as is,
   // preserve attributes.
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     ValidateOpCount<SourceOp, OpCount>();
     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 same operands and result type");
 
     // Cannot convert ops if their operands are not of LLVM type.
     for (Value operand : operands) {
       if (!operand || !operand.getType().isa<LLVM::LLVMType>())
-        return this->matchFailure();
+        return failure();
     }
 
     auto llvmArrayTy = operands[0].getType().cast<LLVM::LLVMType>();
 
     if (!llvmArrayTy.isArrayTy()) {
       auto newOp = rewriter.create<TargetOp>(
           op->getLoc(), operands[0].getType(), operands, op->getAttrs());
       rewriter.replaceOp(op, newOp.getResult());
-      return this->matchSuccess();
+      return success();
     }
 
     if (succeeded(HandleMultidimensionalVectors(
             op, operands, this->typeConverter,
             [&](LLVM::LLVMType llvmVectorTy, ValueRange operands) {
               return rewriter.create<TargetOp>(op->getLoc(), llvmVectorTy,
                                                operands, op->getAttrs());
             },
             rewriter)))
-      return this->matchSuccess();
-    return this->matchFailure();
+      return success();
+    return failure();
   }
 };
 
 template <typename SourceOp, typename TargetOp>
 using UnaryOpLLVMOpLowering = NaryOpLLVMOpLowering<SourceOp, TargetOp, 1>;
 template <typename SourceOp, typename TargetOp>
 using BinaryOpLLVMOpLowering = NaryOpLLVMOpLowering<SourceOp, TargetOp, 2>;
 

 struct AllocOpLowering : public LLVMLegalizationPattern<AllocOp> {
   using LLVMLegalizationPattern<AllocOp>::LLVMLegalizationPattern;
 
   AllocOpLowering(LLVM::LLVMDialect &dialect_, LLVMTypeConverter &converter,
                   bool useAlloca = false)
       : LLVMLegalizationPattern<AllocOp>(dialect_, converter),
         useAlloca(useAlloca) {}
 
-  PatternMatchResult match(Operation *op) const override {
+  LogicalResult match(Operation *op) const override {
     MemRefType type = cast<AllocOp>(op).getType();
     if (isSupportedMemRefType(type))
-      return matchSuccess();
+      return success();
 
     int64_t offset;
     SmallVector<int64_t, 4> strides;
     auto successStrides = getStridesAndOffset(type, strides, offset);
     if (failed(successStrides))
-      return matchFailure();
+      return failure();
 
     // Dynamic strides are ok if they can be deduced from dynamic sizes (which
     // is guaranteed when succeeded(successStrides)). Dynamic offset however can
     // never be alloc'ed.
     if (offset == MemRefType::getDynamicStrideOrOffset())
-      return matchFailure();
+      return failure();
 
-    return matchSuccess();
+    return success();
   }
 
   void rewrite(Operation *op, ArrayRef<Value> operands,
                ConversionPatternRewriter &rewriter) const override {
     auto loc = op->getLoc();
     auto allocOp = cast<AllocOp>(op);
     MemRefType type = allocOp.getType();
 

 // A CallOp automatically promotes MemRefType to a sequence of alloca/store and
 // passes the pointer to the MemRef across function boundaries.
 template <typename CallOpType>
 struct CallOpInterfaceLowering : public LLVMLegalizationPattern<CallOpType> {
   using LLVMLegalizationPattern<CallOpType>::LLVMLegalizationPattern;
   using Super = CallOpInterfaceLowering<CallOpType>;
   using Base = LLVMLegalizationPattern<CallOpType>;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     OperandAdaptor<CallOpType> transformed(operands);
     auto callOp = cast<CallOpType>(op);
 
     // Pack the result types into a struct.
     Type packedResult;
     unsigned numResults = callOp.getNumResults();
     auto resultTypes = llvm::to_vector<4>(callOp.getResultTypes());
 
     for (Type resType : resultTypes) {
       assert(!resType.isa<UnrankedMemRefType>() &&
              "Returning unranked memref is not supported. Pass result as an"
              "argument instead.");
       (void)resType;
     }
 
     if (numResults != 0) {
       if (!(packedResult =
                 this->typeConverter.packFunctionResults(resultTypes)))
-        return this->matchFailure();
+        return failure();
     }
 
     auto promoted = this->typeConverter.promoteMemRefDescriptors(
         op->getLoc(), /*opOperands=*/op->getOperands(), operands, rewriter);
     auto newOp = rewriter.create<LLVM::CallOp>(op->getLoc(), packedResult,
                                                promoted, op->getAttrs());
 
     // If < 2 results, packing did not do anything and we can just return.
     if (numResults < 2) {
       rewriter.replaceOp(op, newOp.getResults());
-      return this->matchSuccess();
+      return success();
     }
 
     // Otherwise, it had been converted to an operation producing a structure.
     // Extract individual results from the structure and return them as list.
     // TODO(aminim, ntv, riverriddle, zinenko): this seems like patching around
     // a particular interaction between MemRefType and CallOp lowering. Find a
     // way to avoid special casing.
     SmallVector<Value, 4> results;
     results.reserve(numResults);
     for (unsigned i = 0; i < numResults; ++i) {
       auto type = this->typeConverter.convertType(op->getResult(i).getType());
       results.push_back(rewriter.create<LLVM::ExtractValueOp>(
           op->getLoc(), type, newOp.getOperation()->getResult(0),
           rewriter.getI64ArrayAttr(i)));
     }
     rewriter.replaceOp(op, results);
 
-    return this->matchSuccess();
+    return success();
   }
 };
 
 struct CallOpLowering : public CallOpInterfaceLowering<CallOp> {
   using Super::Super;
 };
 
 struct CallIndirectOpLowering : public CallOpInterfaceLowering<CallIndirectOp> {
   using Super::Super;
 };
 
 // A `dealloc` is converted into a call to `free` on the underlying data buffer.
 // The memref descriptor being an SSA value, there is no need to clean it up
 // in any way.
 struct DeallocOpLowering : public LLVMLegalizationPattern<DeallocOp> {
   using LLVMLegalizationPattern<DeallocOp>::LLVMLegalizationPattern;
 
   DeallocOpLowering(LLVM::LLVMDialect &dialect_, LLVMTypeConverter &converter,
                     bool useAlloca = false)
       : LLVMLegalizationPattern<DeallocOp>(dialect_, converter),
         useAlloca(useAlloca) {}
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     if (useAlloca)
-      return rewriter.eraseOp(op), matchSuccess();
+      return rewriter.eraseOp(op), success();
 
     assert(operands.size() == 1 && "dealloc takes one operand");
     OperandAdaptor<DeallocOp> transformed(operands);
 
     // Insert the `free` declaration if it is not already present.
     auto freeFunc =
         op->getParentOfType<ModuleOp>().lookupSymbol<LLVM::LLVMFuncOp>("free");
     if (!freeFunc) {
       OpBuilder moduleBuilder(op->getParentOfType<ModuleOp>().getBodyRegion());
       freeFunc = moduleBuilder.create<LLVM::LLVMFuncOp>(
           rewriter.getUnknownLoc(), "free",
           LLVM::LLVMType::getFunctionTy(getVoidType(), getVoidPtrType(),
                                         /*isVarArg=*/false));
     }
 
     MemRefDescriptor memref(transformed.memref());
     Value casted = rewriter.create<LLVM::BitcastOp>(
         op->getLoc(), getVoidPtrType(),
         memref.allocatedPtr(rewriter, op->getLoc()));
     rewriter.replaceOpWithNewOp<LLVM::CallOp>(
         op, ArrayRef<Type>(), rewriter.getSymbolRefAttr(freeFunc), casted);
-    return matchSuccess();
+    return success();
   }
 
   bool useAlloca;
 };
 
 // A `rsqrt` is converted into `1 / sqrt`.
 struct RsqrtOpLowering : public LLVMLegalizationPattern<RsqrtOp> {
   using LLVMLegalizationPattern<RsqrtOp>::LLVMLegalizationPattern;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     OperandAdaptor<RsqrtOp> transformed(operands);
     auto operandType =
         transformed.operand().getType().dyn_cast<LLVM::LLVMType>();
 
     if (!operandType)
-      return matchFailure();
+      return failure();
 
     auto loc = op->getLoc();
     auto resultType = *op->result_type_begin();
     auto floatType = getElementTypeOrSelf(resultType).cast<FloatType>();
     auto floatOne = rewriter.getFloatAttr(floatType, 1.0);
 
     if (!operandType.isArrayTy()) {
       LLVM::ConstantOp one;
       if (operandType.isVectorTy()) {
         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());
       rewriter.replaceOpWithNewOp<LLVM::FDivOp>(op, operandType, one, sqrt);
-      return this->matchSuccess();
+      return success();
     }
 
     auto vectorType = resultType.dyn_cast<VectorType>();
     if (!vectorType)
-      return this->matchFailure();
+      return failure();
 
     if (succeeded(HandleMultidimensionalVectors(
             op, operands, typeConverter,
             [&](LLVM::LLVMType llvmVectorTy, ValueRange operands) {
               auto splatAttr = SplatElementsAttr::get(
                   mlir::VectorType::get({llvmVectorTy.getUnderlyingType()
                                              ->getVectorNumElements()},
                                         floatType),
                   floatOne);
               auto one = rewriter.create<LLVM::ConstantOp>(loc, llvmVectorTy,
                                                            splatAttr);
               auto sqrt =
                   rewriter.create<LLVM::SqrtOp>(loc, llvmVectorTy, operands[0]);
               return rewriter.create<LLVM::FDivOp>(loc, llvmVectorTy, one,
                                                    sqrt);
             },
             rewriter)))
-      return this->matchSuccess();
-    return this->matchFailure();
+      return success();
+    return failure();
   }
 };
 
 // A `tanh` is converted into a call to the `tanh` function.
 struct TanhOpLowering : public LLVMLegalizationPattern<TanhOp> {
   using LLVMLegalizationPattern<TanhOp>::LLVMLegalizationPattern;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
 
     using LLVMFuncOpT = LLVM::LLVMFuncOp;
     using LLVMTypeT = LLVM::LLVMType;
 
     OperandAdaptor<TanhOp> transformed(operands);
     LLVMTypeT operandType =
         transformed.operand().getType().dyn_cast<LLVM::LLVMType>();
 
     if (!operandType)
-      return matchFailure();
+      return failure();
 
     std::string functionName;
     if (operandType.isFloatTy())
       functionName = "tanhf";
     else if (operandType.isDoubleTy())
       functionName = "tanh";
     else
-      return matchFailure();
+      return failure();
 
     // Get a reference to the tanh function, inserting it if necessary.
     Operation *tanhFunc =
         SymbolTable::lookupNearestSymbolFrom(op, functionName);
 
     LLVMFuncOpT tanhLLVMFunc;
     if (tanhFunc) {
       tanhLLVMFunc = cast<LLVMFuncOpT>(tanhFunc);
     } else {
       PatternRewriter::InsertionGuard insertGuard(rewriter);
       auto module = op->getParentOfType<ModuleOp>();
       rewriter.setInsertionPointToStart(module.getBody());
       tanhLLVMFunc = rewriter.create<LLVMFuncOpT>(
           module.getLoc(), functionName,
           LLVMTypeT::getFunctionTy(operandType, operandType,
                                    /*isVarArg=*/false));
     }
 
     rewriter.replaceOpWithNewOp<LLVM::CallOp>(
         op, operandType, rewriter.getSymbolRefAttr(tanhLLVMFunc),
         transformed.operand());
-    return matchSuccess();
+    return success();
   }
 };
 
 struct MemRefCastOpLowering : public LLVMLegalizationPattern<MemRefCastOp> {
   using LLVMLegalizationPattern<MemRefCastOp>::LLVMLegalizationPattern;
 
-  PatternMatchResult match(Operation *op) const override {
+  LogicalResult match(Operation *op) const override {
     auto memRefCastOp = cast<MemRefCastOp>(op);
     Type srcType = memRefCastOp.getOperand().getType();
     Type dstType = memRefCastOp.getType();
 
     if (srcType.isa<MemRefType>() && dstType.isa<MemRefType>()) {
       MemRefType sourceType =
           memRefCastOp.getOperand().getType().cast<MemRefType>();
       MemRefType targetType = memRefCastOp.getType().cast<MemRefType>();
       return (isSupportedMemRefType(targetType) &&
               isSupportedMemRefType(sourceType))
-                 ? matchSuccess()
-                 : matchFailure();
+                 ? success()
+                 : failure();
     }
 
     // At least one of the operands is unranked type
     assert(srcType.isa<UnrankedMemRefType>() ||
            dstType.isa<UnrankedMemRefType>());
 
     // Unranked to unranked cast is disallowed
     return !(srcType.isa<UnrankedMemRefType>() &&
              dstType.isa<UnrankedMemRefType>())
-               ? matchSuccess()
-               : matchFailure();
+               ? success()
+               : failure();
   }
 
   void rewrite(Operation *op, ArrayRef<Value> operands,
                ConversionPatternRewriter &rewriter) const override {
     auto memRefCastOp = cast<MemRefCastOp>(op);
     OperandAdaptor<MemRefCastOp> transformed(operands);
 
     auto srcType = memRefCastOp.getOperand().getType();
void rewrite(Operation *op, ArrayRef<Value> operands,
     }
   }
 };
 
 struct DialectCastOpLowering
     : public LLVMLegalizationPattern<LLVM::DialectCastOp> {
   using LLVMLegalizationPattern<LLVM::DialectCastOp>::LLVMLegalizationPattern;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto castOp = cast<LLVM::DialectCastOp>(op);
     OperandAdaptor<LLVM::DialectCastOp> transformed(operands);
     if (transformed.in().getType() !=
         typeConverter.convertType(castOp.getType())) {
-      return matchFailure();
+      return failure();
     }
     rewriter.replaceOp(op, transformed.in());
-    return matchSuccess();
+    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 LLVMLegalizationPattern<DimOp> {
   using LLVMLegalizationPattern<DimOp>::LLVMLegalizationPattern;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto dimOp = cast<DimOp>(op);
     OperandAdaptor<DimOp> transformed(operands);
     MemRefType type = dimOp.getOperand().getType().cast<MemRefType>();
 
     auto shape = type.getShape();
     int64_t index = dimOp.getIndex();
     // Extract dynamic size from the memref descriptor.
     if (ShapedType::isDynamic(shape[index]))
       rewriter.replaceOp(op, {MemRefDescriptor(transformed.memrefOrTensor())
                                   .size(rewriter, op->getLoc(), index)});
     else
       // Use constant for static size.
       rewriter.replaceOp(
           op, createIndexConstant(rewriter, op->getLoc(), shape[index]));
-    return matchSuccess();
+    return success();
   }
 };
 
 // Common base for load and store operations on MemRefs.  Restricts the match
 // to supported MemRef types.  Provides functionality to emit code accessing a
 // specific element of the underlying data buffer.
 template <typename Derived>
 struct LoadStoreOpLowering : public LLVMLegalizationPattern<Derived> {
   using LLVMLegalizationPattern<Derived>::LLVMLegalizationPattern;
   using Base = LoadStoreOpLowering<Derived>;
 
-  PatternMatchResult match(Operation *op) const override {
+  LogicalResult match(Operation *op) const override {
     MemRefType type = cast<Derived>(op).getMemRefType();
-    return isSupportedMemRefType(type) ? this->matchSuccess()
-                                       : this->matchFailure();
+    return isSupportedMemRefType(type) ? success() : failure();
   }
 
   // Given subscript indices and array sizes in row-major order,
   //   i_n, i_{n-1}, ..., i_1
   //   s_n, s_{n-1}, ..., s_1
   // obtain a value that corresponds to the linearized subscript
   //   \sum_k i_k * \prod_{j=1}^{k-1} s_j
   // by accumulating the running linearized value.
struct LoadStoreOpLowering : public LLVMLegalizationPattern<Derived> {
   }
 };
 
 // Load operation is lowered to obtaining a pointer to the indexed element
 // and loading it.
 struct LoadOpLowering : public LoadStoreOpLowering<LoadOp> {
   using Base::Base;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto loadOp = cast<LoadOp>(op);
     OperandAdaptor<LoadOp> transformed(operands);
     auto type = loadOp.getMemRefType();
 
     Value dataPtr = getDataPtr(op->getLoc(), type, transformed.memref(),
                                transformed.indices(), rewriter, getModule());
     rewriter.replaceOpWithNewOp<LLVM::LoadOp>(op, dataPtr);
-    return matchSuccess();
+    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<StoreOp> {
   using Base::Base;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto type = cast<StoreOp>(op).getMemRefType();
     OperandAdaptor<StoreOp> transformed(operands);
 
     Value dataPtr = getDataPtr(op->getLoc(), type, transformed.memref(),
                                transformed.indices(), rewriter, getModule());
     rewriter.replaceOpWithNewOp<LLVM::StoreOp>(op, transformed.value(),
                                                dataPtr);
-    return matchSuccess();
+    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<PrefetchOp> {
   using Base::Base;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto prefetchOp = cast<PrefetchOp>(op);
     OperandAdaptor<PrefetchOp> transformed(operands);
     auto type = prefetchOp.getMemRefType();
 
     Value dataPtr = getDataPtr(op->getLoc(), type, transformed.memref(),
                                transformed.indices(), rewriter, getModule());
 
     // Replace with llvm.prefetch.
     auto llvmI32Type = typeConverter.convertType(rewriter.getIntegerType(32));
     auto isWrite = rewriter.create<LLVM::ConstantOp>(
         op->getLoc(), llvmI32Type,
         rewriter.getI32IntegerAttr(prefetchOp.isWrite()));
     auto localityHint = rewriter.create<LLVM::ConstantOp>(
         op->getLoc(), llvmI32Type,
         rewriter.getI32IntegerAttr(prefetchOp.localityHint().getZExtValue()));
     auto isData = rewriter.create<LLVM::ConstantOp>(
         op->getLoc(), llvmI32Type,
         rewriter.getI32IntegerAttr(prefetchOp.isDataCache()));
 
     rewriter.replaceOpWithNewOp<LLVM::Prefetch>(op, dataPtr, isWrite,
                                                 localityHint, isData);
-    return matchSuccess();
+    return success();
   }
 };
 
 // The lowering of index_cast becomes an integer conversion since index becomes
 // an integer.  If the bit width of the source and target integer types is the
 // same, just erase the cast.  If the target type is wider, sign-extend the
 // value, otherwise truncate it.
 struct IndexCastOpLowering : public LLVMLegalizationPattern<IndexCastOp> {
   using LLVMLegalizationPattern<IndexCastOp>::LLVMLegalizationPattern;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     IndexCastOpOperandAdaptor transformed(operands);
     auto indexCastOp = cast<IndexCastOp>(op);
 
     auto targetType =
         this->typeConverter.convertType(indexCastOp.getResult().getType())
             .cast<LLVM::LLVMType>();
     auto sourceType = transformed.in().getType().cast<LLVM::LLVMType>();
     unsigned targetBits = targetType.getUnderlyingType()->getIntegerBitWidth();
     unsigned sourceBits = sourceType.getUnderlyingType()->getIntegerBitWidth();
 
     if (targetBits == sourceBits)
       rewriter.replaceOp(op, transformed.in());
     else if (targetBits < sourceBits)
       rewriter.replaceOpWithNewOp<LLVM::TruncOp>(op, targetType,
                                                  transformed.in());
     else
       rewriter.replaceOpWithNewOp<LLVM::SExtOp>(op, targetType,
                                                 transformed.in());
-    return matchSuccess();
+    return success();
   }
 };
 
 // Convert std.cmp predicate into the LLVM dialect CmpPredicate.  The two
 // enums share the numerical values so just cast.
 template <typename LLVMPredType, typename StdPredType>
 static LLVMPredType convertCmpPredicate(StdPredType pred) {
   return static_cast<LLVMPredType>(pred);
 }
 
 struct CmpIOpLowering : public LLVMLegalizationPattern<CmpIOp> {
   using LLVMLegalizationPattern<CmpIOp>::LLVMLegalizationPattern;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto cmpiOp = cast<CmpIOp>(op);
     CmpIOpOperandAdaptor transformed(operands);
 
     rewriter.replaceOpWithNewOp<LLVM::ICmpOp>(
         op, typeConverter.convertType(cmpiOp.getResult().getType()),
         rewriter.getI64IntegerAttr(static_cast<int64_t>(
             convertCmpPredicate<LLVM::ICmpPredicate>(cmpiOp.getPredicate()))),
         transformed.lhs(), transformed.rhs());
 
-    return matchSuccess();
+    return success();
   }
 };
 
 struct CmpFOpLowering : public LLVMLegalizationPattern<CmpFOp> {
   using LLVMLegalizationPattern<CmpFOp>::LLVMLegalizationPattern;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto cmpfOp = cast<CmpFOp>(op);
     CmpFOpOperandAdaptor transformed(operands);
 
     rewriter.replaceOpWithNewOp<LLVM::FCmpOp>(
         op, typeConverter.convertType(cmpfOp.getResult().getType()),
         rewriter.getI64IntegerAttr(static_cast<int64_t>(
             convertCmpPredicate<LLVM::FCmpPredicate>(cmpfOp.getPredicate()))),
         transformed.lhs(), transformed.rhs());
 
-    return matchSuccess();
+    return success();
   }
 };
 
 struct SIToFPLowering
     : public OneToOneLLVMOpLowering<SIToFPOp, LLVM::SIToFPOp> {
   using Super::Super;
 };
 

 
 // Base class for LLVM IR lowering terminator operations with successors.
 template <typename SourceOp, typename TargetOp>
 struct OneToOneLLVMTerminatorLowering
     : public LLVMLegalizationPattern<SourceOp> {
   using LLVMLegalizationPattern<SourceOp>::LLVMLegalizationPattern;
   using Super = OneToOneLLVMTerminatorLowering<SourceOp, TargetOp>;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     rewriter.replaceOpWithNewOp<TargetOp>(op, operands, op->getSuccessors(),
                                           op->getAttrs());
-    return this->matchSuccess();
+    return success();
   }
 };
 
 // Special lowering pattern for `ReturnOps`.  Unlike all other operations,
 // `ReturnOp` interacts with the function signature and must have as many
 // operands as the function has return values.  Because in LLVM IR, functions
 // can only return 0 or 1 value, we pack multiple values into a structure type.
 // Emit `UndefOp` followed by `InsertValueOp`s to create such structure if
 // necessary before returning it
 struct ReturnOpLowering : public LLVMLegalizationPattern<ReturnOp> {
   using LLVMLegalizationPattern<ReturnOp>::LLVMLegalizationPattern;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     unsigned numArguments = op->getNumOperands();
 
     // If ReturnOp has 0 or 1 operand, create it and return immediately.
     if (numArguments == 0) {
       rewriter.replaceOpWithNewOp<LLVM::ReturnOp>(
           op, ArrayRef<Type>(), ArrayRef<Value>(), op->getAttrs());
-      return matchSuccess();
+      return success();
     }
     if (numArguments == 1) {
       rewriter.replaceOpWithNewOp<LLVM::ReturnOp>(
           op, ArrayRef<Type>(), operands.front(), op->getAttrs());
-      return matchSuccess();
+      return success();
     }
 
     // Otherwise, we need to pack the arguments into an LLVM struct type before
     // returning.
     auto packedType = typeConverter.packFunctionResults(
         llvm::to_vector<4>(op->getOperandTypes()));
 
     Value packed = rewriter.create<LLVM::UndefOp>(op->getLoc(), packedType);
     for (unsigned i = 0; i < numArguments; ++i) {
       packed = rewriter.create<LLVM::InsertValueOp>(
           op->getLoc(), packedType, packed, operands[i],
           rewriter.getI64ArrayAttr(i));
     }
     rewriter.replaceOpWithNewOp<LLVM::ReturnOp>(op, ArrayRef<Type>(), packed,
                                                 op->getAttrs());
-    return matchSuccess();
+    return success();
   }
 };
 
 // FIXME: this should be tablegen'ed as well.
 struct BranchOpLowering
     : public OneToOneLLVMTerminatorLowering<BranchOp, LLVM::BrOp> {
   using Super::Super;
 };
 struct CondBranchOpLowering
     : public OneToOneLLVMTerminatorLowering<CondBranchOp, LLVM::CondBrOp> {
   using Super::Super;
 };
 
 // The Splat operation is lowered to an insertelement + a shufflevector
 // operation. Splat to only 1-d vector result types are lowered.
 struct SplatOpLowering : public LLVMLegalizationPattern<SplatOp> {
   using LLVMLegalizationPattern<SplatOp>::LLVMLegalizationPattern;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto splatOp = cast<SplatOp>(op);
     VectorType resultType = splatOp.getType().dyn_cast<VectorType>();
     if (!resultType || resultType.getRank() != 1)
-      return matchFailure();
+      return failure();
 
     // First insert it into an undef vector so we can shuffle it.
     auto vectorType = typeConverter.convertType(splatOp.getType());
     Value undef = rewriter.create<LLVM::UndefOp>(op->getLoc(), vectorType);
     auto zero = rewriter.create<LLVM::ConstantOp>(
         op->getLoc(), typeConverter.convertType(rewriter.getIntegerType(32)),
         rewriter.getZeroAttr(rewriter.getIntegerType(32)));
 
     auto v = rewriter.create<LLVM::InsertElementOp>(
         op->getLoc(), vectorType, undef, splatOp.getOperand(), zero);
 
     int64_t width = splatOp.getType().cast<VectorType>().getDimSize(0);
     SmallVector<int32_t, 4> zeroValues(width, 0);
 
     // Shuffle the value across the desired number of elements.
     ArrayAttr zeroAttrs = rewriter.getI32ArrayAttr(zeroValues);
     rewriter.replaceOpWithNewOp<LLVM::ShuffleVectorOp>(op, v, undef, zeroAttrs);
-    return matchSuccess();
+    return success();
   }
 };
 
 // The Splat operation is lowered to an insertelement + a shufflevector
 // operation. Splat to only 2+-d vector result types are lowered by the
 // SplatNdOpLowering, the 1-d case is handled by SplatOpLowering.
 struct SplatNdOpLowering : public LLVMLegalizationPattern<SplatOp> {
   using LLVMLegalizationPattern<SplatOp>::LLVMLegalizationPattern;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto splatOp = cast<SplatOp>(op);
     OperandAdaptor<SplatOp> adaptor(operands);
     VectorType resultType = splatOp.getType().dyn_cast<VectorType>();
     if (!resultType || resultType.getRank() == 1)
-      return matchFailure();
+      return failure();
 
     // First insert it into an undef vector so we can shuffle it.
     auto loc = op->getLoc();
     auto vectorTypeInfo = extractNDVectorTypeInfo(resultType, typeConverter);
     auto llvmArrayTy = vectorTypeInfo.llvmArrayTy;
     auto llvmVectorTy = vectorTypeInfo.llvmVectorTy;
     if (!llvmArrayTy || !llvmVectorTy)
-      return matchFailure();
+      return failure();
 
     // Construct returned value.
     Value desc = rewriter.create<LLVM::UndefOp>(loc, llvmArrayTy);
 
     // Construct a 1-D vector with the splatted value that we insert in all the
     // places within the returned descriptor.
     Value vdesc = rewriter.create<LLVM::UndefOp>(loc, llvmVectorTy);
     auto zero = rewriter.create<LLVM::ConstantOp>(
matchAndRewrite(Operation *op, ArrayRef<Value> operands,
 
     // Iterate of linear index, convert to coords space and insert splatted 1-D
     // vector in each position.
     nDVectorIterate(vectorTypeInfo, rewriter, [&](ArrayAttr position) {
       desc = rewriter.create<LLVM::InsertValueOp>(loc, llvmArrayTy, desc, v,
                                                   position);
     });
     rewriter.replaceOp(op, desc);
-    return matchSuccess();
+    return success();
   }
 };
 
 /// Conversion pattern that transforms a subview op into:
 ///   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 LLVMLegalizationPattern<SubViewOp> {
   using LLVMLegalizationPattern<SubViewOp>::LLVMLegalizationPattern;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto loc = op->getLoc();
     auto viewOp = cast<SubViewOp>(op);
     // TODO(b/144779634, ravishankarm) : After Tblgen is adapted to support
     // having multiple variadic operands where each operand can have different
     // number of entries, clean all of this up.
     SmallVector<Value, 2> dynamicOffsets(
matchAndRewrite(Operation *op, ArrayRef<Value> operands,
 
     auto viewMemRefType = viewOp.getType();
     auto targetElementTy =
         typeConverter.convertType(viewMemRefType.getElementType())
             .dyn_cast<LLVM::LLVMType>();
     auto targetDescTy = typeConverter.convertType(viewMemRefType)
                             .dyn_cast_or_null<LLVM::LLVMType>();
     if (!sourceElementTy || !targetDescTy)
-      return matchFailure();
+      return failure();
 
     // Currently, only rank > 0 and full or no operands are supported. Fail to
     // convert otherwise.
     unsigned rank = sourceMemRefType.getRank();
     if (viewMemRefType.getRank() == 0 ||
         (!dynamicOffsets.empty() && rank != dynamicOffsets.size()) ||
         (!dynamicSizes.empty() && rank != dynamicSizes.size()) ||
         (!dynamicStrides.empty() && rank != dynamicStrides.size()))
-      return matchFailure();
+      return failure();
 
     int64_t offset;
     SmallVector<int64_t, 4> strides;
     auto successStrides = getStridesAndOffset(viewMemRefType, strides, offset);
     if (failed(successStrides))
-      return matchFailure();
+      return failure();
 
     // Fail to convert if neither a dynamic nor static offset is available.
     if (dynamicOffsets.empty() &&
         offset == MemRefType::getDynamicStrideOrOffset())
-      return matchFailure();
+      return failure();
 
     // Create the descriptor.
     if (!operands.front().getType().isa<LLVM::LLVMType>())
-      return matchFailure();
+      return failure();
     MemRefDescriptor sourceMemRef(operands.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, targetElementTy.getPointerTo(), extracted);
     targetMemRef.setAllocatedPtr(rewriter, loc, bitcastPtr);
for (int i = viewMemRefType.getRank() - 1; i >= 0; --i) {
             loc, llvmIndexType, rewriter.getI64IntegerAttr(strides[i]));
       else
         newStride = rewriter.create<LLVM::MulOp>(loc, dynamicStrides[i],
                                                  strideValues[i]);
       targetMemRef.setStride(rewriter, loc, i, newStride);
     }
 
     rewriter.replaceOp(op, {targetMemRef});
-    return matchSuccess();
+    return success();
   }
 };
 
 /// Conversion pattern that transforms a op into:
 ///   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 view op is replaced by the descriptor.
Value getStride(ConversionPatternRewriter &rewriter, Location loc,
     if (nextSize)
       return runningStride
                  ? rewriter.create<LLVM::MulOp>(loc, runningStride, nextSize)
                  : nextSize;
     assert(!runningStride);
     return createIndexConstant(rewriter, loc, 1);
   }
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto loc = op->getLoc();
     auto viewOp = cast<ViewOp>(op);
     ViewOpOperandAdaptor adaptor(operands);
 
     auto viewMemRefType = viewOp.getType();
     auto targetElementTy =
         typeConverter.convertType(viewMemRefType.getElementType())
             .dyn_cast<LLVM::LLVMType>();
     auto targetDescTy =
         typeConverter.convertType(viewMemRefType).dyn_cast<LLVM::LLVMType>();
     if (!targetDescTy)
       return op->emitWarning("Target descriptor type not converted to LLVM"),
-             matchFailure();
+             failure();
 
     int64_t offset;
     SmallVector<int64_t, 4> strides;
     auto successStrides = getStridesAndOffset(viewMemRefType, strides, offset);
     if (failed(successStrides))
-      return op->emitWarning("cannot cast to non-strided shape"),
-             matchFailure();
+      return op->emitWarning("cannot cast to non-strided shape"), failure();
 
     // Create the descriptor.
     MemRefDescriptor sourceMemRef(adaptor.source());
     auto targetMemRef = MemRefDescriptor::undef(rewriter, loc, targetDescTy);
 
     // Field 1: Copy the allocated pointer, used for malloc/free.
     Value extracted = sourceMemRef.allocatedPtr(rewriter, loc);
     Value bitcastPtr = rewriter.create<LLVM::BitcastOp>(
matchAndRewrite(Operation *op, ArrayRef<Value> operands,
     Value baseOffset = !hasDynamicOffset
                            ? createIndexConstant(rewriter, loc, offset)
                            // TODO(ntv): better adaptor.
                            : sizeAndOffsetOperands.front();
     targetMemRef.setOffset(rewriter, loc, baseOffset);
 
     // Early exit for 0-D corner case.
     if (viewMemRefType.getRank() == 0)
-      return rewriter.replaceOp(op, {targetMemRef}), matchSuccess();
+      return rewriter.replaceOp(op, {targetMemRef}), success();
 
     // Fields 4 and 5: Update sizes and strides.
     if (strides.back() != 1)
-      return op->emitWarning("cannot cast to non-contiguous shape"),
-             matchFailure();
+      return op->emitWarning("cannot cast to non-contiguous shape"), failure();
     Value stride = nullptr, nextSize = nullptr;
     // Drop the dynamic stride from the operand list, if present.
     ArrayRef<Value> sizeOperands(sizeAndOffsetOperands);
     if (hasDynamicOffset)
       sizeOperands = sizeOperands.drop_front();
     for (int i = viewMemRefType.getRank() - 1; i >= 0; --i) {
       // Update size.
       Value size =
           getSize(rewriter, loc, viewMemRefType.getShape(), sizeOperands, i);
       targetMemRef.setSize(rewriter, loc, i, size);
       // Update stride.
       stride = getStride(rewriter, loc, strides, nextSize, stride, i);
       targetMemRef.setStride(rewriter, loc, i, stride);
       nextSize = size;
     }
 
     rewriter.replaceOp(op, {targetMemRef});
-    return matchSuccess();
+    return success();
   }
 };
 
 struct AssumeAlignmentOpLowering
     : public LLVMLegalizationPattern<AssumeAlignmentOp> {
   using LLVMLegalizationPattern<AssumeAlignmentOp>::LLVMLegalizationPattern;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     OperandAdaptor<AssumeAlignmentOp> transformed(operands);
     Value memref = transformed.memref();
     unsigned alignment = cast<AssumeAlignmentOp>(op).alignment().getZExtValue();
 
     MemRefDescriptor memRefDescriptor(memref);
     Value ptr = memRefDescriptor.alignedPtr(rewriter, memref.getLoc());
Value ptrValue =
         rewriter.create<LLVM::PtrToIntOp>(op->getLoc(), getIndexType(), ptr);
     rewriter.create<LLVM::AssumeOp>(
         op->getLoc(),
         rewriter.create<LLVM::ICmpOp>(
             op->getLoc(), LLVM::ICmpPredicate::eq,
             rewriter.create<LLVM::AndOp>(op->getLoc(), ptrValue, mask), zero));
 
     rewriter.eraseOp(op);
-    return matchSuccess();
+    return success();
   }
 };
 
 } // namespace
 
 /// Try to match the kind of a std.atomic_rmw to determine whether to use a
 /// lowering to llvm.atomicrmw or fallback to llvm.cmpxchg.
 static Optional<LLVM::AtomicBinOp> matchSimpleAtomicOp(AtomicRMWOp atomicOp) {
static Optional<LLVM::AtomicBinOp> matchSimpleAtomicOp(AtomicRMWOp atomicOp) {
   llvm_unreachable("Invalid AtomicRMWKind");
 }
 
 namespace {
 
 struct AtomicRMWOpLowering : public LoadStoreOpLowering<AtomicRMWOp> {
   using Base::Base;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto atomicOp = cast<AtomicRMWOp>(op);
     auto maybeKind = matchSimpleAtomicOp(atomicOp);
     if (!maybeKind)
-      return matchFailure();
+      return failure();
     OperandAdaptor<AtomicRMWOp> adaptor(operands);
     auto resultType = adaptor.value().getType();
     auto memRefType = atomicOp.getMemRefType();
     auto dataPtr = getDataPtr(op->getLoc(), memRefType, adaptor.memref(),
                               adaptor.indices(), rewriter, getModule());
     rewriter.replaceOpWithNewOp<LLVM::AtomicRMWOp>(
         op, resultType, *maybeKind, dataPtr, adaptor.value(),
         LLVM::AtomicOrdering::acq_rel);
-    return matchSuccess();
+    return success();
   }
 };
 
 /// Wrap a llvm.cmpxchg operation in a while loop so that the operation can be
 /// retried until it succeeds in atomically storing a new value into memory.
 ///
 ///      +---------------------------------+
 ///      |   <code before the AtomicRMWOp> |

 ///      +--------------------------------+
 ///      | end:                           |
 ///      |   <code after the AtomicRMWOp> |
 ///      +--------------------------------+
 ///
 struct AtomicCmpXchgOpLowering : public LoadStoreOpLowering<AtomicRMWOp> {
   using Base::Base;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto atomicOp = cast<AtomicRMWOp>(op);
     auto maybeKind = matchSimpleAtomicOp(atomicOp);
     if (maybeKind)
-      return matchFailure();
+      return failure();
 
     LLVM::FCmpPredicate predicate;
     switch (atomicOp.kind()) {
     case AtomicRMWKind::maxf:
       predicate = LLVM::FCmpPredicate::ogt;
       break;
     case AtomicRMWKind::minf:
       predicate = LLVM::FCmpPredicate::olt;
       break;
     default:
-      return matchFailure();
+      return failure();
     }
 
     OperandAdaptor<AtomicRMWOp> adaptor(operands);
     auto loc = op->getLoc();
     auto valueType = adaptor.value().getType().cast<LLVM::LLVMType>();
 
     // Split the block into initial, loop, and ending parts.
     auto *initBlock = rewriter.getInsertionBlock();
matchAndRewrite(Operation *op, ArrayRef<Value> operands,
 
     // Conditionally branch to the end or back to the loop depending on %ok.
     rewriter.create<LLVM::CondBrOp>(loc, ok, endBlock, ArrayRef<Value>(),
                                     loopBlock, newLoaded);
 
     // The 'result' of the atomic_rmw op is the newly loaded value.
     rewriter.replaceOp(op, {newLoaded});
 
-    return matchSuccess();
+    return success();
   }
 };
 
 } // namespace
 
 static void ensureDistinctSuccessors(Block &bb) {
   Operation *terminator = bb.getTerminator();
   if (terminator->getNumSuccessors() < 2)

mlir/lib/Conversion/StandardToSPIRV/ConvertStandardToSPIRV.cpp

 namespace {
 
 /// Convert composite constant operation to SPIR-V dialect.
 // TODO(denis0x0D) : move to DRR.
 class ConstantCompositeOpConversion final : public SPIRVOpLowering<ConstantOp> {
 public:
   using SPIRVOpLowering<ConstantOp>::SPIRVOpLowering;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(ConstantOp constCompositeOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
 /// Convert constant operation with IndexType return to SPIR-V constant
 /// operation. Since IndexType is not used within SPIR-V dialect, this needs
 /// special handling to make sure the result type and the type of the value
 /// attribute are consistent.
 // TODO(ravishankarm) : This should be moved into DRR.
 class ConstantIndexOpConversion final : public SPIRVOpLowering<ConstantOp> {
 public:
   using SPIRVOpLowering<ConstantOp>::SPIRVOpLowering;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(ConstantOp constIndexOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
 /// Convert floating-point comparison operations to SPIR-V dialect.
 class CmpFOpConversion final : public SPIRVOpLowering<CmpFOp> {
 public:
   using SPIRVOpLowering<CmpFOp>::SPIRVOpLowering;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(CmpFOp cmpFOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
 /// Convert compare operation to SPIR-V dialect.
 class CmpIOpConversion final : public SPIRVOpLowering<CmpIOp> {
 public:
   using SPIRVOpLowering<CmpIOp>::SPIRVOpLowering;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(CmpIOp cmpIOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
 /// Convert integer binary operations to SPIR-V operations. Cannot use
 /// tablegen for this. If the integer operation is on variables of IndexType,
 /// the type of the return value of the replacement operation differs from
 /// that of the replaced operation. This is not handled in tablegen-based
 /// pattern specification.
 // TODO(ravishankarm) : This should be moved into DRR.
 template <typename StdOp, typename SPIRVOp>
 class IntegerOpConversion final : public SPIRVOpLowering<StdOp> {
 public:
   using SPIRVOpLowering<StdOp>::SPIRVOpLowering;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(StdOp operation, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto resultType =
         this->typeConverter.convertType(operation.getResult().getType());
     rewriter.template replaceOpWithNewOp<SPIRVOp>(
         operation, resultType, operands, ArrayRef<NamedAttribute>());
-    return this->matchSuccess();
+    return success();
   }
 };
 
 /// Convert load -> spv.LoadOp. The operands of the replaced operation are of
 /// IndexType while that of the replacement operation are of type i32. This is
 /// not supported in tablegen based pattern specification.
 // TODO(ravishankarm) : This should be moved into DRR.
 class LoadOpConversion final : public SPIRVOpLowering<LoadOp> {
 public:
   using SPIRVOpLowering<LoadOp>::SPIRVOpLowering;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(LoadOp loadOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
 /// Convert return -> spv.Return.
 // TODO(ravishankarm) : This should be moved into DRR.
 class ReturnOpConversion final : public SPIRVOpLowering<ReturnOp> {
 public:
   using SPIRVOpLowering<ReturnOp>::SPIRVOpLowering;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(ReturnOp returnOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
 /// Convert select -> spv.Select
 // TODO(ravishankarm) : This should be moved into DRR.
 class SelectOpConversion final : public SPIRVOpLowering<SelectOp> {
 public:
   using SPIRVOpLowering<SelectOp>::SPIRVOpLowering;
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(SelectOp op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
 /// Convert store -> spv.StoreOp. The operands of the replaced operation are
 /// of IndexType while that of the replacement operation are of type i32. This
 /// is not supported in tablegen based pattern specification.
 // TODO(ravishankarm) : This should be moved into DRR.
 class StoreOpConversion final : public SPIRVOpLowering<StoreOp> {
 public:
   using SPIRVOpLowering<StoreOp>::SPIRVOpLowering;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(StoreOp storeOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
 } // namespace
 
 //===----------------------------------------------------------------------===//
 // ConstantOp with composite type.
 //===----------------------------------------------------------------------===//
 
-PatternMatchResult ConstantCompositeOpConversion::matchAndRewrite(
+LogicalResult ConstantCompositeOpConversion::matchAndRewrite(
     ConstantOp constCompositeOp, ArrayRef<Value> operands,
     ConversionPatternRewriter &rewriter) const {
   auto compositeType =
       constCompositeOp.getResult().getType().dyn_cast<RankedTensorType>();
   if (!compositeType)
-    return matchFailure();
+    return failure();
 
   auto spirvCompositeType = typeConverter.convertType(compositeType);
   if (!spirvCompositeType)
-    return matchFailure();
+    return failure();
 
   auto linearizedElements =
       constCompositeOp.value().dyn_cast<DenseElementsAttr>();
   if (!linearizedElements)
-    return matchFailure();
+    return failure();
 
   // If composite type has rank greater than one, then perform linearization.
   if (compositeType.getRank() > 1) {
     auto linearizedType = RankedTensorType::get(compositeType.getNumElements(),
                                                 compositeType.getElementType());
     linearizedElements = linearizedElements.reshape(linearizedType);
   }
 
   rewriter.replaceOpWithNewOp<spirv::ConstantOp>(
       constCompositeOp, spirvCompositeType, linearizedElements);
-  return matchSuccess();
+  return success();
 }
 
 //===----------------------------------------------------------------------===//
 // ConstantOp with index type.
 //===----------------------------------------------------------------------===//
 
-PatternMatchResult ConstantIndexOpConversion::matchAndRewrite(
+LogicalResult ConstantIndexOpConversion::matchAndRewrite(
     ConstantOp constIndexOp, ArrayRef<Value> operands,
     ConversionPatternRewriter &rewriter) const {
   if (!constIndexOp.getResult().getType().isa<IndexType>()) {
-    return matchFailure();
+    return failure();
   }
   // The attribute has index type which is not directly supported in
   // SPIR-V. Get the integer value and create a new IntegerAttr.
   auto constAttr = constIndexOp.value().dyn_cast<IntegerAttr>();
   if (!constAttr) {
-    return matchFailure();
+    return failure();
   }
 
   // Use the bitwidth set in the value attribute to decide the result type
   // of the SPIR-V constant operation since SPIR-V does not support index
   // types.
   auto constVal = constAttr.getValue();
   auto constValType = constAttr.getType().dyn_cast<IndexType>();
   if (!constValType) {
-    return matchFailure();
+    return failure();
   }
   auto spirvConstType =
       typeConverter.convertType(constIndexOp.getResult().getType());
   auto spirvConstVal =
       rewriter.getIntegerAttr(spirvConstType, constAttr.getInt());
   rewriter.replaceOpWithNewOp<spirv::ConstantOp>(constIndexOp, spirvConstType,
                                                  spirvConstVal);
-  return matchSuccess();
+  return success();
 }
 
 //===----------------------------------------------------------------------===//
 // CmpFOp
 //===----------------------------------------------------------------------===//
 
-PatternMatchResult
+LogicalResult
 CmpFOpConversion::matchAndRewrite(CmpFOp cmpFOp, ArrayRef<Value> operands,
                                   ConversionPatternRewriter &rewriter) const {
   CmpFOpOperandAdaptor cmpFOpOperands(operands);
 
   switch (cmpFOp.getPredicate()) {
 #define DISPATCH(cmpPredicate, spirvOp)                                        \
   case cmpPredicate:                                                           \
     rewriter.replaceOpWithNewOp<spirvOp>(cmpFOp, cmpFOp.getResult().getType(), \
                                          cmpFOpOperands.lhs(),                 \
                                          cmpFOpOperands.rhs());                \
-    return matchSuccess();
+    return success();
 
     // Ordered.
     DISPATCH(CmpFPredicate::OEQ, spirv::FOrdEqualOp);
     DISPATCH(CmpFPredicate::OGT, spirv::FOrdGreaterThanOp);
     DISPATCH(CmpFPredicate::OGE, spirv::FOrdGreaterThanEqualOp);
     DISPATCH(CmpFPredicate::OLT, spirv::FOrdLessThanOp);
     DISPATCH(CmpFPredicate::OLE, spirv::FOrdLessThanEqualOp);
     DISPATCH(CmpFPredicate::ONE, spirv::FOrdNotEqualOp);
     // Unordered.
     DISPATCH(CmpFPredicate::UEQ, spirv::FUnordEqualOp);
     DISPATCH(CmpFPredicate::UGT, spirv::FUnordGreaterThanOp);
     DISPATCH(CmpFPredicate::UGE, spirv::FUnordGreaterThanEqualOp);
     DISPATCH(CmpFPredicate::ULT, spirv::FUnordLessThanOp);
     DISPATCH(CmpFPredicate::ULE, spirv::FUnordLessThanEqualOp);
     DISPATCH(CmpFPredicate::UNE, spirv::FUnordNotEqualOp);
 
 #undef DISPATCH
 
   default:
     break;
   }
-  return matchFailure();
+  return failure();
 }
 
 //===----------------------------------------------------------------------===//
 // CmpIOp
 //===----------------------------------------------------------------------===//
 
-PatternMatchResult
+LogicalResult
 CmpIOpConversion::matchAndRewrite(CmpIOp cmpIOp, ArrayRef<Value> operands,
                                   ConversionPatternRewriter &rewriter) const {
   CmpIOpOperandAdaptor cmpIOpOperands(operands);
 
   switch (cmpIOp.getPredicate()) {
 #define DISPATCH(cmpPredicate, spirvOp)                                        \
   case cmpPredicate:                                                           \
     rewriter.replaceOpWithNewOp<spirvOp>(cmpIOp, cmpIOp.getResult().getType(), \
                                          cmpIOpOperands.lhs(),                 \
                                          cmpIOpOperands.rhs());                \
-    return matchSuccess();
+    return success();
 
     DISPATCH(CmpIPredicate::eq, spirv::IEqualOp);
     DISPATCH(CmpIPredicate::ne, spirv::INotEqualOp);
     DISPATCH(CmpIPredicate::slt, spirv::SLessThanOp);
     DISPATCH(CmpIPredicate::sle, spirv::SLessThanEqualOp);
     DISPATCH(CmpIPredicate::sgt, spirv::SGreaterThanOp);
     DISPATCH(CmpIPredicate::sge, spirv::SGreaterThanEqualOp);
     DISPATCH(CmpIPredicate::ult, spirv::ULessThanOp);
     DISPATCH(CmpIPredicate::ule, spirv::ULessThanEqualOp);
     DISPATCH(CmpIPredicate::ugt, spirv::UGreaterThanOp);
     DISPATCH(CmpIPredicate::uge, spirv::UGreaterThanEqualOp);
 
 #undef DISPATCH
   }
-  return matchFailure();
+  return failure();
 }
 
 //===----------------------------------------------------------------------===//
 // LoadOp
 //===----------------------------------------------------------------------===//
 
-PatternMatchResult
+LogicalResult
 LoadOpConversion::matchAndRewrite(LoadOp loadOp, ArrayRef<Value> operands,
                                   ConversionPatternRewriter &rewriter) const {
   LoadOpOperandAdaptor loadOperands(operands);
   auto loadPtr = spirv::getElementPtr(
       typeConverter, loadOp.memref().getType().cast<MemRefType>(),
       loadOperands.memref(), loadOperands.indices(), loadOp.getLoc(), rewriter);
   rewriter.replaceOpWithNewOp<spirv::LoadOp>(loadOp, loadPtr);
-  return matchSuccess();
+  return success();
 }
 
 //===----------------------------------------------------------------------===//
 // ReturnOp
 //===----------------------------------------------------------------------===//
 
-PatternMatchResult
+LogicalResult
 ReturnOpConversion::matchAndRewrite(ReturnOp returnOp, ArrayRef<Value> operands,
                                     ConversionPatternRewriter &rewriter) const {
   if (returnOp.getNumOperands()) {
-    return matchFailure();
+    return failure();
   }
   rewriter.replaceOpWithNewOp<spirv::ReturnOp>(returnOp);
-  return matchSuccess();
+  return success();
 }
 
 //===----------------------------------------------------------------------===//
 // SelectOp
 //===----------------------------------------------------------------------===//
 
-PatternMatchResult
+LogicalResult
 SelectOpConversion::matchAndRewrite(SelectOp op, ArrayRef<Value> operands,
                                     ConversionPatternRewriter &rewriter) const {
   SelectOpOperandAdaptor selectOperands(operands);
   rewriter.replaceOpWithNewOp<spirv::SelectOp>(op, selectOperands.condition(),
                                                selectOperands.true_value(),
                                                selectOperands.false_value());
-  return matchSuccess();
+  return success();
 }
 
 //===----------------------------------------------------------------------===//
 // StoreOp
 //===----------------------------------------------------------------------===//
 
-PatternMatchResult
+LogicalResult
 StoreOpConversion::matchAndRewrite(StoreOp storeOp, ArrayRef<Value> operands,
                                    ConversionPatternRewriter &rewriter) const {
   StoreOpOperandAdaptor storeOperands(operands);
   auto storePtr = spirv::getElementPtr(
       typeConverter, storeOp.memref().getType().cast<MemRefType>(),
       storeOperands.memref(), storeOperands.indices(), storeOp.getLoc(),
       rewriter);
   rewriter.replaceOpWithNewOp<spirv::StoreOp>(storeOp, storePtr,
                                               storeOperands.value());
-  return matchSuccess();
+  return success();
 }
 
 namespace {
 /// Import the Standard Ops to SPIR-V Patterns.
 #include "StandardToSPIRV.cpp.inc"
 } // namespace
 
 namespace mlir {

mlir/lib/Conversion/StandardToSPIRV/LegalizeStandardForSPIRV.cpp

 using namespace mlir;
 
 namespace {
 /// Merges subview operation with load operation.
 class LoadOpOfSubViewFolder final : public OpRewritePattern<LoadOp> {
 public:
   using OpRewritePattern<LoadOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(LoadOp loadOp,
-                                     PatternRewriter &rewriter) const override;
+  LogicalResult matchAndRewrite(LoadOp loadOp,
 };
 
 /// Merges subview operation with store operation.
 class StoreOpOfSubViewFolder final : public OpRewritePattern<StoreOp> {
 public:
   using OpRewritePattern<StoreOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(StoreOp storeOp,
-                                     PatternRewriter &rewriter) const override;
+  LogicalResult matchAndRewrite(StoreOp storeOp,
 };
 } // namespace
 
 //===----------------------------------------------------------------------===//
 // Utility functions for op legalization.
 //===----------------------------------------------------------------------===//
 
 /// Given the 'indices' of an load/store operation where the memref is a result
resolveSourceIndices(Location loc, PatternRewriter &rewriter,
   }
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // Folding SubViewOp and LoadOp.
 //===----------------------------------------------------------------------===//
 
-PatternMatchResult
+LogicalResult
 LoadOpOfSubViewFolder::matchAndRewrite(LoadOp loadOp,
                                        PatternRewriter &rewriter) const {
   auto subViewOp = dyn_cast_or_null<SubViewOp>(loadOp.memref().getDefiningOp());
   if (!subViewOp) {
-    return matchFailure();
+    return failure();
   }
   SmallVector<Value, 4> sourceIndices;
   if (failed(resolveSourceIndices(loadOp.getLoc(), rewriter, subViewOp,
                                   loadOp.indices(), sourceIndices)))
-    return matchFailure();
+    return failure();
 
   rewriter.replaceOpWithNewOp<LoadOp>(loadOp, subViewOp.source(),
                                       sourceIndices);
-  return matchSuccess();
+  return success();
 }
 
 //===----------------------------------------------------------------------===//
 // Folding SubViewOp and StoreOp.
 //===----------------------------------------------------------------------===//
 
-PatternMatchResult
+LogicalResult
 StoreOpOfSubViewFolder::matchAndRewrite(StoreOp storeOp,
                                         PatternRewriter &rewriter) const {
   auto subViewOp =
       dyn_cast_or_null<SubViewOp>(storeOp.memref().getDefiningOp());
   if (!subViewOp) {
-    return matchFailure();
+    return failure();
   }
   SmallVector<Value, 4> sourceIndices;
   if (failed(resolveSourceIndices(storeOp.getLoc(), rewriter, subViewOp,
                                   storeOp.indices(), sourceIndices)))
-    return matchFailure();
+    return failure();
 
   rewriter.replaceOpWithNewOp<StoreOp>(storeOp, storeOp.value(),
                                        subViewOp.source(), sourceIndices);
-  return matchSuccess();
+  return success();
 }
 
 //===----------------------------------------------------------------------===//
 // Hook for adding patterns.
 //===----------------------------------------------------------------------===//
 
 void mlir::populateStdLegalizationPatternsForSPIRVLowering(
     MLIRContext *context, OwningRewritePatternList &patterns) {

mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVM.cpp

 
 class VectorBroadcastOpConversion : public ConvertToLLVMPattern {
 public:
   explicit VectorBroadcastOpConversion(MLIRContext *context,
                                        LLVMTypeConverter &typeConverter)
       : ConvertToLLVMPattern(vector::BroadcastOp::getOperationName(), context,
                              typeConverter) {}
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto broadcastOp = cast<vector::BroadcastOp>(op);
     VectorType dstVectorType = broadcastOp.getVectorType();
     if (typeConverter.convertType(dstVectorType) == nullptr)
-      return matchFailure();
+      return failure();
     // Rewrite when the full vector type can be lowered (which
     // implies all 'reduced' types can be lowered too).
     auto adaptor = vector::BroadcastOpOperandAdaptor(operands);
     VectorType srcVectorType =
         broadcastOp.getSourceType().dyn_cast<VectorType>();
     rewriter.replaceOp(
         op, expandRanks(adaptor.source(), // source value to be expanded
                         op->getLoc(),     // location of original broadcast
                         srcVectorType, dstVectorType, rewriter));
-    return matchSuccess();
+    return success();
   }
 
 private:
   // Expands the given source value over all the ranks, as defined
   // by the source and destination type (a null source type denotes
   // expansion from a scalar value into a vector).
   //
   // TODO(ajcbik): consider replacing this one-pattern lowering

 /// This is lowered directly to the proper llvm.intr.matrix.multiply.
 class VectorMatmulOpConversion : public ConvertToLLVMPattern {
 public:
   explicit VectorMatmulOpConversion(MLIRContext *context,
                                     LLVMTypeConverter &typeConverter)
       : ConvertToLLVMPattern(vector::MatmulOp::getOperationName(), context,
                              typeConverter) {}
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto matmulOp = cast<vector::MatmulOp>(op);
     auto adaptor = vector::MatmulOpOperandAdaptor(operands);
     rewriter.replaceOpWithNewOp<LLVM::MatrixMultiplyOp>(
         op, typeConverter.convertType(matmulOp.res().getType()), adaptor.lhs(),
         adaptor.rhs(), matmulOp.lhs_rows(), matmulOp.lhs_columns(),
         matmulOp.rhs_columns());
-    return matchSuccess();
+    return success();
   }
 };
 
 class VectorReductionOpConversion : public ConvertToLLVMPattern {
 public:
   explicit VectorReductionOpConversion(MLIRContext *context,
                                        LLVMTypeConverter &typeConverter)
       : ConvertToLLVMPattern(vector::ReductionOp::getOperationName(), context,
                              typeConverter) {}
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto reductionOp = cast<vector::ReductionOp>(op);
     auto kind = reductionOp.kind();
     Type eltType = reductionOp.dest().getType();
     Type llvmType = typeConverter.convertType(eltType);
     if (eltType.isSignlessInteger(32) || eltType.isSignlessInteger(64)) {
       // Integer reductions: add/mul/min/max/and/or/xor.
if (eltType.isSignlessInteger(32) || eltType.isSignlessInteger(64)) {
             op, llvmType, operands[0]);
       else if (kind == "or")
         rewriter.replaceOpWithNewOp<LLVM::experimental_vector_reduce_or>(
             op, llvmType, operands[0]);
       else if (kind == "xor")
         rewriter.replaceOpWithNewOp<LLVM::experimental_vector_reduce_xor>(
             op, llvmType, operands[0]);
       else
-        return matchFailure();
-      return matchSuccess();
+        return failure();
+      return success();
 
     } else if (eltType.isF32() || eltType.isF64()) {
       // Floating-point reductions: add/mul/min/max
       if (kind == "add") {
         // Optional accumulator (or zero).
         Value acc = operands.size() > 1 ? operands[1]
                                         : rewriter.create<LLVM::ConstantOp>(
                                               op->getLoc(), llvmType,
if (eltType.isSignlessInteger(32) || eltType.isSignlessInteger(64)) {
             op, llvmType, acc, operands[0]);
       } else if (kind == "min")
         rewriter.replaceOpWithNewOp<LLVM::experimental_vector_reduce_fmin>(
             op, llvmType, operands[0]);
       else if (kind == "max")
         rewriter.replaceOpWithNewOp<LLVM::experimental_vector_reduce_fmax>(
             op, llvmType, operands[0]);
       else
-        return matchFailure();
-      return matchSuccess();
+        return failure();
+      return success();
     }
-    return matchFailure();
+    return failure();
   }
 };
 
 class VectorShuffleOpConversion : public ConvertToLLVMPattern {
 public:
   explicit VectorShuffleOpConversion(MLIRContext *context,
                                      LLVMTypeConverter &typeConverter)
       : ConvertToLLVMPattern(vector::ShuffleOp::getOperationName(), context,
                              typeConverter) {}
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto loc = op->getLoc();
     auto adaptor = vector::ShuffleOpOperandAdaptor(operands);
     auto shuffleOp = cast<vector::ShuffleOp>(op);
     auto v1Type = shuffleOp.getV1VectorType();
     auto v2Type = shuffleOp.getV2VectorType();
     auto vectorType = shuffleOp.getVectorType();
     Type llvmType = typeConverter.convertType(vectorType);
     auto maskArrayAttr = shuffleOp.mask();
 
     // Bail if result type cannot be lowered.
     if (!llvmType)
-      return matchFailure();
+      return failure();
 
     // Get rank and dimension sizes.
     int64_t rank = vectorType.getRank();
     assert(v1Type.getRank() == rank);
     assert(v2Type.getRank() == rank);
     int64_t v1Dim = v1Type.getDimSize(0);
 
     // For rank 1, where both operands have *exactly* the same vector type,
     // there is direct shuffle support in LLVM. Use it!
     if (rank == 1 && v1Type == v2Type) {
       Value shuffle = rewriter.create<LLVM::ShuffleVectorOp>(
           loc, adaptor.v1(), adaptor.v2(), maskArrayAttr);
       rewriter.replaceOp(op, shuffle);
-      return matchSuccess();
+      return success();
     }
 
     // For all other cases, insert the individual values individually.
     Value insert = rewriter.create<LLVM::UndefOp>(loc, llvmType);
     int64_t insPos = 0;
     for (auto en : llvm::enumerate(maskArrayAttr)) {
       int64_t extPos = en.value().cast<IntegerAttr>().getInt();
       Value value = adaptor.v1();
       if (extPos >= v1Dim) {
         extPos -= v1Dim;
         value = adaptor.v2();
       }
       Value extract = extractOne(rewriter, typeConverter, loc, value, llvmType,
                                  rank, extPos);
       insert = insertOne(rewriter, typeConverter, loc, insert, extract,
                          llvmType, rank, insPos++);
     }
     rewriter.replaceOp(op, insert);
-    return matchSuccess();
+    return success();
   }
 };
 
 class VectorExtractElementOpConversion : public ConvertToLLVMPattern {
 public:
   explicit VectorExtractElementOpConversion(MLIRContext *context,
                                             LLVMTypeConverter &typeConverter)
       : ConvertToLLVMPattern(vector::ExtractElementOp::getOperationName(),
                              context, typeConverter) {}
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto adaptor = vector::ExtractElementOpOperandAdaptor(operands);
     auto extractEltOp = cast<vector::ExtractElementOp>(op);
     auto vectorType = extractEltOp.getVectorType();
     auto llvmType = typeConverter.convertType(vectorType.getElementType());
 
     // Bail if result type cannot be lowered.
     if (!llvmType)
-      return matchFailure();
+      return failure();
 
     rewriter.replaceOpWithNewOp<LLVM::ExtractElementOp>(
         op, llvmType, adaptor.vector(), adaptor.position());
-    return matchSuccess();
+    return success();
   }
 };
 
 class VectorExtractOpConversion : public ConvertToLLVMPattern {
 public:
   explicit VectorExtractOpConversion(MLIRContext *context,
                                      LLVMTypeConverter &typeConverter)
       : ConvertToLLVMPattern(vector::ExtractOp::getOperationName(), context,
                              typeConverter) {}
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto loc = op->getLoc();
     auto adaptor = vector::ExtractOpOperandAdaptor(operands);
     auto extractOp = cast<vector::ExtractOp>(op);
     auto vectorType = extractOp.getVectorType();
     auto resultType = extractOp.getResult().getType();
     auto llvmResultType = typeConverter.convertType(resultType);
     auto positionArrayAttr = extractOp.position();
 
     // Bail if result type cannot be lowered.
     if (!llvmResultType)
-      return matchFailure();
+      return failure();
 
     // One-shot extraction of vector from array (only requires extractvalue).
     if (resultType.isa<VectorType>()) {
       Value extracted = rewriter.create<LLVM::ExtractValueOp>(
           loc, llvmResultType, adaptor.vector(), positionArrayAttr);
       rewriter.replaceOp(op, extracted);
-      return matchSuccess();
+      return success();
     }
 
     // Potential extraction of 1-D vector from array.
     auto *context = op->getContext();
     Value extracted = adaptor.vector();
     auto positionAttrs = positionArrayAttr.getValue();
     if (positionAttrs.size() > 1) {
       auto oneDVectorType = reducedVectorTypeBack(vectorType);
       auto nMinusOnePositionAttrs =
           ArrayAttr::get(positionAttrs.drop_back(), context);
       extracted = rewriter.create<LLVM::ExtractValueOp>(
           loc, typeConverter.convertType(oneDVectorType), extracted,
           nMinusOnePositionAttrs);
     }
 
     // Remaining extraction of element from 1-D LLVM vector
     auto position = positionAttrs.back().cast<IntegerAttr>();
     auto i64Type = LLVM::LLVMType::getInt64Ty(typeConverter.getDialect());
     auto constant = rewriter.create<LLVM::ConstantOp>(loc, i64Type, position);
     extracted =
         rewriter.create<LLVM::ExtractElementOp>(loc, extracted, constant);
     rewriter.replaceOp(op, extracted);
 
-    return matchSuccess();
+    return success();
   }
 };
 
 /// Conversion pattern that turns a vector.fma on a 1-D vector
 /// into an llvm.intr.fmuladd. This is a trivial 1-1 conversion.
 /// This does not match vectors of n >= 2 rank.
 ///
 /// Example:
 /// ```
 ///  vector.fma %a, %a, %a : vector<8xf32>
 /// ```
 /// is converted to:
 /// ```
 ///  llvm.intr.fma %va, %va, %va:
 ///    (!llvm<"<8 x float>">, !llvm<"<8 x float>">, !llvm<"<8 x float>">)
 ///    -> !llvm<"<8 x float>">
 /// ```
 class VectorFMAOp1DConversion : public ConvertToLLVMPattern {
 public:
   explicit VectorFMAOp1DConversion(MLIRContext *context,
                                    LLVMTypeConverter &typeConverter)
       : ConvertToLLVMPattern(vector::FMAOp::getOperationName(), context,
                              typeConverter) {}
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto adaptor = vector::FMAOpOperandAdaptor(operands);
     vector::FMAOp fmaOp = cast<vector::FMAOp>(op);
     VectorType vType = fmaOp.getVectorType();
     if (vType.getRank() != 1)
-      return matchFailure();
+      return failure();
     rewriter.replaceOpWithNewOp<LLVM::FMAOp>(op, adaptor.lhs(), adaptor.rhs(),
                                              adaptor.acc());
-    return matchSuccess();
+    return success();
   }
 };
 
 class VectorInsertElementOpConversion : public ConvertToLLVMPattern {
 public:
   explicit VectorInsertElementOpConversion(MLIRContext *context,
                                            LLVMTypeConverter &typeConverter)
       : ConvertToLLVMPattern(vector::InsertElementOp::getOperationName(),
                              context, typeConverter) {}
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto adaptor = vector::InsertElementOpOperandAdaptor(operands);
     auto insertEltOp = cast<vector::InsertElementOp>(op);
     auto vectorType = insertEltOp.getDestVectorType();
     auto llvmType = typeConverter.convertType(vectorType);
 
     // Bail if result type cannot be lowered.
     if (!llvmType)
-      return matchFailure();
+      return failure();
 
     rewriter.replaceOpWithNewOp<LLVM::InsertElementOp>(
         op, llvmType, adaptor.dest(), adaptor.source(), adaptor.position());
-    return matchSuccess();
+    return success();
   }
 };
 
 class VectorInsertOpConversion : public ConvertToLLVMPattern {
 public:
   explicit VectorInsertOpConversion(MLIRContext *context,
                                     LLVMTypeConverter &typeConverter)
       : ConvertToLLVMPattern(vector::InsertOp::getOperationName(), context,
                              typeConverter) {}
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto loc = op->getLoc();
     auto adaptor = vector::InsertOpOperandAdaptor(operands);
     auto insertOp = cast<vector::InsertOp>(op);
     auto sourceType = insertOp.getSourceType();
     auto destVectorType = insertOp.getDestVectorType();
     auto llvmResultType = typeConverter.convertType(destVectorType);
     auto positionArrayAttr = insertOp.position();
 
     // Bail if result type cannot be lowered.
     if (!llvmResultType)
-      return matchFailure();
+      return failure();
 
     // One-shot insertion of a vector into an array (only requires insertvalue).
     if (sourceType.isa<VectorType>()) {
       Value inserted = rewriter.create<LLVM::InsertValueOp>(
           loc, llvmResultType, adaptor.dest(), adaptor.source(),
           positionArrayAttr);
       rewriter.replaceOp(op, inserted);
-      return matchSuccess();
+      return success();
     }
 
     // Potential extraction of 1-D vector from array.
     auto *context = op->getContext();
     Value extracted = adaptor.dest();
     auto positionAttrs = positionArrayAttr.getValue();
     auto position = positionAttrs.back().cast<IntegerAttr>();
     auto oneDVectorType = destVectorType;
if (positionAttrs.size() > 1) {
       auto nMinusOnePositionAttrs =
           ArrayAttr::get(positionAttrs.drop_back(), context);
       inserted = rewriter.create<LLVM::InsertValueOp>(loc, llvmResultType,
                                                       adaptor.dest(), inserted,
                                                       nMinusOnePositionAttrs);
     }
 
     rewriter.replaceOp(op, inserted);
-    return matchSuccess();
+    return success();
   }
 };
 
 /// Rank reducing rewrite for n-D FMA into (n-1)-D FMA where n > 1.
 ///
 /// Example:
 /// ```
 ///   %d = vector.fma %a, %b, %c : vector<2x4xf32>

 ///  %vd2 = vector.fma %va2, %vb2, %vc2 : vector<4xf32>
 ///  %r3 = vector.insertvalue %vd2, %r2[1] : vector<4xf32> into vector<2x4xf32>
 ///  // %r3 holds the final value.
 /// ```
 class VectorFMAOpNDRewritePattern : public OpRewritePattern<FMAOp> {
 public:
   using OpRewritePattern<FMAOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(FMAOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(FMAOp op,
     auto vType = op.getVectorType();
     if (vType.getRank() < 2)
-      return matchFailure();
+      return failure();
 
     auto loc = op.getLoc();
     auto elemType = vType.getElementType();
     Value zero = rewriter.create<ConstantOp>(loc, elemType,
                                              rewriter.getZeroAttr(elemType));
     Value desc = rewriter.create<SplatOp>(loc, vType, zero);
     for (int64_t i = 0, e = vType.getShape().front(); i != e; ++i) {
       Value extrLHS = rewriter.create<ExtractOp>(loc, op.lhs(), i);
       Value extrRHS = rewriter.create<ExtractOp>(loc, op.rhs(), i);
       Value extrACC = rewriter.create<ExtractOp>(loc, op.acc(), i);
       Value fma = rewriter.create<FMAOp>(loc, extrLHS, extrRHS, extrACC);
       desc = rewriter.create<InsertOp>(loc, fma, desc, i);
     }
     rewriter.replaceOp(op, desc);
-    return matchSuccess();
+    return success();
   }
 };
 
 // When ranks are different, InsertStridedSlice needs to extract a properly
 // ranked vector from the destination vector into which to insert. This pattern
 // only takes care of this part and forwards the rest of the conversion to
 // another pattern that converts InsertStridedSlice for operands of the same
 // rank.
 //
 // RewritePattern for InsertStridedSliceOp where source and destination vectors
 // have different ranks. In this case:
 //   1. the proper subvector is extracted from the destination vector
 //   2. a new InsertStridedSlice op is created to insert the source in the
 //   destination subvector
 //   3. the destination subvector is inserted back in the proper place
 //   4. the op is replaced by the result of step 3.
 // The new InsertStridedSlice from step 2. will be picked up by a
 // `VectorInsertStridedSliceOpSameRankRewritePattern`.
 class VectorInsertStridedSliceOpDifferentRankRewritePattern
     : public OpRewritePattern<InsertStridedSliceOp> {
 public:
   using OpRewritePattern<InsertStridedSliceOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(InsertStridedSliceOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(InsertStridedSliceOp op,
     auto srcType = op.getSourceVectorType();
     auto dstType = op.getDestVectorType();
 
     if (op.offsets().getValue().empty())
-      return matchFailure();
+      return failure();
 
     auto loc = op.getLoc();
     int64_t rankDiff = dstType.getRank() - srcType.getRank();
     assert(rankDiff >= 0);
     if (rankDiff == 0)
-      return matchFailure();
+      return failure();
 
     int64_t rankRest = dstType.getRank() - rankDiff;
     // Extract / insert the subvector of matching rank and InsertStridedSlice
     // on it.
     Value extracted =
         rewriter.create<ExtractOp>(loc, op.dest(),
                                    getI64SubArray(op.offsets(), /*dropFront=*/0,
                                                   /*dropFront=*/rankRest));
     // A different pattern will kick in for InsertStridedSlice with matching
     // ranks.
     auto stridedSliceInnerOp = rewriter.create<InsertStridedSliceOp>(
         loc, op.source(), extracted,
         getI64SubArray(op.offsets(), /*dropFront=*/rankDiff),
         getI64SubArray(op.strides(), /*dropFront=*/0));
     rewriter.replaceOpWithNewOp<InsertOp>(
         op, stridedSliceInnerOp.getResult(), op.dest(),
         getI64SubArray(op.offsets(), /*dropFront=*/0,
                        /*dropFront=*/rankRest));
-    return matchSuccess();
+    return success();
   }
 };
 
 // RewritePattern for InsertStridedSliceOp where source and destination vectors
 // have the same rank. In this case, we reduce
 //   1. the proper subvector is extracted from the destination vector
 //   2. a new InsertStridedSlice op is created to insert the source in the
 //   destination subvector
 //   3. the destination subvector is inserted back in the proper place
 //   4. the op is replaced by the result of step 3.
 // The new InsertStridedSlice from step 2. will be picked up by a
 // `VectorInsertStridedSliceOpSameRankRewritePattern`.
 class VectorInsertStridedSliceOpSameRankRewritePattern
     : public OpRewritePattern<InsertStridedSliceOp> {
 public:
   using OpRewritePattern<InsertStridedSliceOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(InsertStridedSliceOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(InsertStridedSliceOp op,
     auto srcType = op.getSourceVectorType();
     auto dstType = op.getDestVectorType();
 
     if (op.offsets().getValue().empty())
-      return matchFailure();
+      return failure();
 
     int64_t rankDiff = dstType.getRank() - srcType.getRank();
     assert(rankDiff >= 0);
     if (rankDiff != 0)
-      return matchFailure();
+      return failure();
 
     if (srcType == dstType) {
       rewriter.replaceOp(op, op.source());
-      return matchSuccess();
+      return success();
     }
 
     int64_t offset =
         op.offsets().getValue().front().cast<IntegerAttr>().getInt();
     int64_t size = srcType.getShape().front();
     int64_t stride =
         op.strides().getValue().front().cast<IntegerAttr>().getInt();
 
for (int64_t off = offset, e = offset + size * stride, idx = 0; off < e;
         assert(succeeded(success) && "Unexpected failure");
         extractedSource = insertStridedSliceOp;
       }
       // 4. Insert the extractedSource into the res vector.
       res = insertOne(rewriter, loc, extractedSource, res, off);
     }
 
     rewriter.replaceOp(op, res);
-    return matchSuccess();
+    return success();
   }
 };
 
 class VectorTypeCastOpConversion : public ConvertToLLVMPattern {
 public:
   explicit VectorTypeCastOpConversion(MLIRContext *context,
                                       LLVMTypeConverter &typeConverter)
       : ConvertToLLVMPattern(vector::TypeCastOp::getOperationName(), context,
                              typeConverter) {}
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto loc = op->getLoc();
     vector::TypeCastOp castOp = cast<vector::TypeCastOp>(op);
     MemRefType sourceMemRefType =
         castOp.getOperand().getType().cast<MemRefType>();
     MemRefType targetMemRefType =
         castOp.getResult().getType().cast<MemRefType>();
 
     // Only static shape casts supported atm.
     if (!sourceMemRefType.hasStaticShape() ||
         !targetMemRefType.hasStaticShape())
-      return matchFailure();
+      return failure();
 
     auto llvmSourceDescriptorTy =
         operands[0].getType().dyn_cast<LLVM::LLVMType>();
     if (!llvmSourceDescriptorTy || !llvmSourceDescriptorTy.isStructTy())
-      return matchFailure();
+      return failure();
     MemRefDescriptor sourceMemRef(operands[0]);
 
     auto llvmTargetDescriptorTy = typeConverter.convertType(targetMemRefType)
                                       .dyn_cast_or_null<LLVM::LLVMType>();
     if (!llvmTargetDescriptorTy || !llvmTargetDescriptorTy.isStructTy())
-      return matchFailure();
+      return failure();
 
     int64_t offset;
     SmallVector<int64_t, 4> strides;
     auto successStrides =
         getStridesAndOffset(sourceMemRefType, strides, offset);
     bool isContiguous = (strides.back() == 1);
     if (isContiguous) {
       auto sizes = sourceMemRefType.getShape();
       for (int index = 0, e = strides.size() - 2; index < e; ++index) {
         if (strides[index] != strides[index + 1] * sizes[index + 1]) {
           isContiguous = false;
           break;
         }
       }
     }
     // Only contiguous source tensors supported atm.
     if (failed(successStrides) || !isContiguous)
-      return matchFailure();
+      return failure();
 
     auto int64Ty = LLVM::LLVMType::getInt64Ty(typeConverter.getDialect());
 
     // Create descriptor.
     auto desc = MemRefDescriptor::undef(rewriter, loc, llvmTargetDescriptorTy);
     Type llvmTargetElementTy = desc.getElementType();
     // Set allocated ptr.
     Value allocated = sourceMemRef.allocatedPtr(rewriter, loc);
for (auto indexedSize : llvm::enumerate(targetMemRefType.getShape())) {
       desc.setSize(rewriter, loc, index, size);
       auto strideAttr =
           rewriter.getIntegerAttr(rewriter.getIndexType(), strides[index]);
       auto stride = rewriter.create<LLVM::ConstantOp>(loc, int64Ty, strideAttr);
       desc.setStride(rewriter, loc, index, stride);
     }
 
     rewriter.replaceOp(op, {desc});
-    return matchSuccess();
+    return success();
   }
 };
 
 class VectorPrintOpConversion : public ConvertToLLVMPattern {
 public:
   explicit VectorPrintOpConversion(MLIRContext *context,
                                    LLVMTypeConverter &typeConverter)
       : ConvertToLLVMPattern(vector::PrintOp::getOperationName(), context,
                              typeConverter) {}
 
   // Proof-of-concept lowering implementation that relies on a small
   // runtime support library, which only needs to provide a few
   // printing methods (single value for all data types, opening/closing
   // bracket, comma, newline). The lowering fully unrolls a vector
   // in terms of these elementary printing operations. The advantage
   // of this approach is that the library can remain unaware of all
   // low-level implementation details of vectors while still supporting
   // output of any shaped and dimensioned vector. Due to full unrolling,
   // this approach is less suited for very large vectors though.
   //
   // TODO(ajcbik): rely solely on libc in future? something else?
   //
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto printOp = cast<vector::PrintOp>(op);
     auto adaptor = vector::PrintOpOperandAdaptor(operands);
     Type printType = printOp.getPrintType();
 
     if (typeConverter.convertType(printType) == nullptr)
-      return matchFailure();
+      return failure();
 
     // Make sure element type has runtime support (currently just Float/Double).
     VectorType vectorType = printType.dyn_cast<VectorType>();
     Type eltType = vectorType ? vectorType.getElementType() : printType;
     int64_t rank = vectorType ? vectorType.getRank() : 0;
     Operation *printer;
     if (eltType.isSignlessInteger(32))
       printer = getPrintI32(op);
     else if (eltType.isSignlessInteger(64))
       printer = getPrintI64(op);
     else if (eltType.isF32())
       printer = getPrintFloat(op);
     else if (eltType.isF64())
       printer = getPrintDouble(op);
     else
-      return matchFailure();
+      return failure();
 
     // Unroll vector into elementary print calls.
     emitRanks(rewriter, op, adaptor.source(), vectorType, printer, rank);
     emitCall(rewriter, op->getLoc(), getPrintNewline(op));
     rewriter.eraseOp(op);
-    return matchSuccess();
+    return success();
   }
 
 private:
   void emitRanks(ConversionPatternRewriter &rewriter, Operation *op,
                  Value value, VectorType vectorType, Operation *printer,
                  int64_t rank) const {
     Location loc = op->getLoc();
     if (rank == 0) {

 
 /// Progressive lowering of StridedSliceOp to either:
 ///   1. extractelement + insertelement for the 1-D case
 ///   2. extract + optional strided_slice + insert for the n-D case.
 class VectorStridedSliceOpConversion : public OpRewritePattern<StridedSliceOp> {
 public:
   using OpRewritePattern<StridedSliceOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(StridedSliceOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(StridedSliceOp op,
     auto dstType = op.getResult().getType().cast<VectorType>();
 
     assert(!op.offsets().getValue().empty() && "Unexpected empty offsets");
 
     int64_t offset =
         op.offsets().getValue().front().cast<IntegerAttr>().getInt();
     int64_t size = op.sizes().getValue().front().cast<IntegerAttr>().getInt();
     int64_t stride =
for (int64_t off = offset, e = offset + size * stride, idx = 0; off < e;
         auto success = matchAndRewrite(stridedSliceOp, rewriter);
         (void)success;
         assert(succeeded(success) && "Unexpected failure");
         extracted = stridedSliceOp;
       }
       res = insertOne(rewriter, loc, extracted, res, idx);
     }
     rewriter.replaceOp(op, {res});
-    return matchSuccess();
+    return success();
   }
 };
 
 } // namespace
 
 /// Populate the given list with patterns that convert from Vector to LLVM.
 void mlir::populateVectorToLLVMConversionPatterns(
     LLVMTypeConverter &converter, OwningRewritePatternList &patterns) {

mlir/lib/Conversion/VectorToLoops/ConvertVectorToLoops.cpp

struct VectorTransferRewriter : public RewritePattern {
   /// Used for staging the transfer in a local scalar buffer.
   MemRefType tmpMemRefType(TransferOpTy transfer) const {
     auto vectorType = transfer.getVectorType();
     return MemRefType::get(vectorType.getShape(), vectorType.getElementType(),
                            {}, 0);
   }
 
   /// Performs the rewrite.
-  PatternMatchResult matchAndRewrite(Operation *op,
-                                     PatternRewriter &rewriter) const override;
+  LogicalResult matchAndRewrite(Operation *op,
 };
 
 /// Lowers TransferReadOp into a combination of:
 ///   1. local memory allocation;
 ///   2. perfect loop nest over:
 ///      a. scalar load from local buffers (viewed as a scalar memref);
 ///      a. scalar store to original memref (with clipping).
 ///   3. vector_load from local buffer (viewed as a memref<1 x vector>);

 /// in the presence of data-parallel only operations, we generate code that
 /// writes the same value multiple time on the edge locations.
 ///
 /// TODO(ntv): implement alternatives to clipping.
 /// TODO(ntv): support non-data-parallel operations.
 
 /// Performs the rewrite.
 template <>
-PatternMatchResult VectorTransferRewriter<TransferReadOp>::matchAndRewrite(
+LogicalResult VectorTransferRewriter<TransferReadOp>::matchAndRewrite(
     Operation *op, PatternRewriter &rewriter) const {
   using namespace mlir::edsc::op;
 
   TransferReadOp transfer = cast<TransferReadOp>(op);
 
   // 1. Setup all the captures.
   ScopedContext scope(rewriter, transfer.getLoc());
   StdIndexedValue remote(transfer.memref());
LoopNestBuilder(pivs, lbs, ubs, steps)([&] {
     // Computes clippedScalarAccessExprs in the loop nest scope (ivs exist).
     local(ivs) = remote(clip(transfer, memRefBoundsCapture, ivs));
   });
   ValueHandle vectorValue = std_load(vec);
   (std_dealloc(tmp)); // vexing parse
 
   // 3. Propagate.
   rewriter.replaceOp(op, vectorValue.getValue());
-  return matchSuccess();
+  return success();
 }
 
 /// Lowers TransferWriteOp into a combination of:
 ///   1. local memory allocation;
 ///   2. vector_store to local buffer (viewed as a memref<1 x vector>);
 ///   3. perfect loop nest over:
 ///      a. scalar load from local buffers (viewed as a scalar memref);
 ///      a. scalar store to original memref (with clipping).
 ///   4. local memory deallocation.
 ///
 /// More specifically, lowers the data transfer part while ensuring no
 /// out-of-bounds accesses are possible. Out-of-bounds behavior is handled by
 /// clipping. This means that a given value in memory can be written to multiple
 /// times and concurrently.
 ///
 /// See `Important notes about clipping and full-tiles only abstraction` in the
 /// description of `readClipped` above.
 ///
 /// TODO(ntv): implement alternatives to clipping.
 /// TODO(ntv): support non-data-parallel operations.
 template <>
-PatternMatchResult VectorTransferRewriter<TransferWriteOp>::matchAndRewrite(
+LogicalResult VectorTransferRewriter<TransferWriteOp>::matchAndRewrite(
     Operation *op, PatternRewriter &rewriter) const {
   using namespace edsc::op;
 
   TransferWriteOp transfer = cast<TransferWriteOp>(op);
 
   // 1. Setup all the captures.
   ScopedContext scope(rewriter, transfer.getLoc());
   StdIndexedValue remote(transfer.memref());
LogicalResult VectorTransferRewriter<TransferWriteOp>::matchAndRewrite(
   std_store(vectorValue, vec);
   LoopNestBuilder(pivs, lbs, ubs, steps)([&] {
     // Computes clippedScalarAccessExprs in the loop nest scope (ivs exist).
     remote(clip(transfer, memRefBoundsCapture, ivs)) = local(ivs);
   });
   (std_dealloc(tmp)); // vexing parse...
 
   rewriter.eraseOp(op);
-  return matchSuccess();
+  return success();
 }
 
 } // namespace
 
 void mlir::populateVectorToAffineLoopsConversionPatterns(
     MLIRContext *context, OwningRewritePatternList &patterns) {
   patterns.insert<VectorTransferRewriter<vector::TransferReadOp>,
                   VectorTransferRewriter<vector::TransferWriteOp>>(context);
 }

mlir/lib/Dialect/AffineOps/AffineOps.cpp

 struct SimplifyAffineOp : public OpRewritePattern<AffineOpTy> {
   using OpRewritePattern<AffineOpTy>::OpRewritePattern;
 
   /// Replace the affine op with another instance of it with the supplied
   /// map and mapOperands.
   void replaceAffineOp(PatternRewriter &rewriter, AffineOpTy affineOp,
                        AffineMap map, ArrayRef<Value> mapOperands) const;
 
-  PatternMatchResult matchAndRewrite(AffineOpTy affineOp,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(AffineOpTy affineOp,
     static_assert(std::is_same<AffineOpTy, AffineLoadOp>::value ||
                       std::is_same<AffineOpTy, AffinePrefetchOp>::value ||
                       std::is_same<AffineOpTy, AffineStoreOp>::value ||
                       std::is_same<AffineOpTy, AffineApplyOp>::value ||
                       std::is_same<AffineOpTy, AffineMinOp>::value ||
                       std::is_same<AffineOpTy, AffineMaxOp>::value,
                   "affine load/store/apply op expected");
     auto map = affineOp.getAffineMap();
     AffineMap oldMap = map;
     auto oldOperands = affineOp.getMapOperands();
     SmallVector<Value, 8> resultOperands(oldOperands);
     composeAffineMapAndOperands(&map, &resultOperands);
     if (map == oldMap && std::equal(oldOperands.begin(), oldOperands.end(),
                                     resultOperands.begin()))
-      return this->matchFailure();
+      return failure();
 
     replaceAffineOp(rewriter, affineOp, map, resultOperands);
-    return this->matchSuccess();
+    return success();
   }
 };
 
 // Specialize the template to account for the different build signatures for
 // affine load, store, and apply ops.
 template <>
 void SimplifyAffineOp<AffineLoadOp>::replaceAffineOp(
     PatternRewriter &rewriter, AffineLoadOp load, AffineMap map,
static LogicalResult canonicalizeLoopBounds(AffineForOp forOp) {
   return success();
 }
 
 namespace {
 /// This is a pattern to fold trivially empty loops.
 struct AffineForEmptyLoopFolder : public OpRewritePattern<AffineForOp> {
   using OpRewritePattern<AffineForOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(AffineForOp forOp,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(AffineForOp forOp,
     // Check that the body only contains a terminator.
     if (!has_single_element(*forOp.getBody()))
-      return matchFailure();
+      return failure();
     rewriter.eraseOp(forOp);
-    return matchSuccess();
+    return success();
   }
 };
 } // end anonymous namespace
 
 void AffineForOp::getCanonicalizationPatterns(OwningRewritePatternList &results,
                                               MLIRContext *context) {
   results.insert<AffineForEmptyLoopFolder>(context);
 }

mlir/lib/Dialect/FxpMathOps/Transforms/LowerUniformRealMath.cpp

static Value emitDequantize(Location loc, Value input,
   }
 }
 
 namespace {
 
 struct UniformDequantizePattern : public OpRewritePattern<DequantizeCastOp> {
   using OpRewritePattern<DequantizeCastOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(DequantizeCastOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(DequantizeCastOp op,
     Type inputType = op.arg().getType();
     Type outputType = op.getResult().getType();
 
     QuantizedType inputElementType =
         QuantizedType::getQuantizedElementType(inputType);
     Type expressedOutputType = inputElementType.castToExpressedType(inputType);
     if (expressedOutputType != outputType) {
       // Not a valid uniform cast.
-      return matchFailure();
+      return failure();
     }
 
     Value dequantizedValue = emitDequantize(op.getLoc(), op.arg(), rewriter);
     if (!dequantizedValue) {
-      return matchFailure();
+      return failure();
     }
 
     rewriter.replaceOp(op, dequantizedValue);
-    return matchSuccess();
+    return success();
   }
 };
 
 } // end anonymous namespace
 
 //===----------------------------------------------------------------------===//
 // Elementwise add
 //===----------------------------------------------------------------------===//
tryRewriteAffineMulEwSigned(const UniformBinaryOpInfo &info,
   return success();
 }
 
 namespace {
 
 struct UniformRealAddEwPattern : public OpRewritePattern<RealAddEwOp> {
   using OpRewritePattern<RealAddEwOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(RealAddEwOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(RealAddEwOp op,
     const UniformBinaryOpInfo info(op, op.lhs(), op.rhs(), op.clamp_min(),
                                    op.clamp_max());
     if (!info.isValid()) {
-      return matchFailure();
+      return failure();
     }
 
     // Try all of the permutations we support.
     if (succeeded(tryRewriteAffineAddEwIsomorphicSigned(info, rewriter))) {
-      return matchSuccess();
+      return success();
     }
 
-    return matchFailure();
+    return failure();
   }
 };
 
 struct UniformRealMulEwPattern : public OpRewritePattern<RealMulEwOp> {
   using OpRewritePattern<RealMulEwOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(RealMulEwOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(RealMulEwOp op,
     const UniformBinaryOpInfo info(op, op.lhs(), op.rhs(), op.clamp_min(),
                                    op.clamp_max());
     if (!info.isValid()) {
-      return matchFailure();
+      return failure();
     }
 
     // Try all of the permutations we support.
     if (succeeded(tryRewriteAffineMulEwSigned(info, rewriter))) {
-      return matchSuccess();
+      return success();
     }
 
-    return matchFailure();
+    return failure();
   }
 };
 
 } // end anonymous namespace
 
 //===----------------------------------------------------------------------===//
 // LowerUniformRealMath pass
 //===----------------------------------------------------------------------===//

mlir/lib/Dialect/GPU/Transforms/AllReduceLowering.cpp

private:
 
   static constexpr int kSubgroupSize = 32;
 };
 
 struct GpuAllReduceConversion : public RewritePattern {
   explicit GpuAllReduceConversion(MLIRContext *context)
       : RewritePattern(gpu::GPUFuncOp::getOperationName(), 1, context) {}
 
-  PatternMatchResult matchAndRewrite(Operation *op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(Operation *op,
     auto funcOp = cast<gpu::GPUFuncOp>(op);
     auto callback = [&](gpu::AllReduceOp reduceOp) {
       GpuAllReduceRewriter(funcOp, reduceOp, rewriter).rewrite();
       // Performing a rewrite invalidates the walk iterator. Report interrupt
       // so that we can start a new walk until all all_reduce ops are replaced.
       return WalkResult::interrupt();
     };
     while (funcOp.walk(callback).wasInterrupted()) {
     }
-    return matchSuccess();
+    return success();
   }
 };
 } // namespace
 
 void mlir::populateGpuRewritePatterns(MLIRContext *context,
                                       OwningRewritePatternList &patterns) {
   patterns.insert<GpuAllReduceConversion>(context);
 }

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

 }
 
 namespace {
 
 /// Patterns to fuse a generic op, with the producer of its operands.
 struct FuseGenericTensorOps : public OpRewritePattern<GenericOp> {
   using OpRewritePattern<GenericOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(GenericOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(GenericOp op,
     if (!op.hasTensorSemantics())
-      return matchFailure();
+      return failure();
 
     // Find the first operand that is defined by another generic op on tensors.
     for (auto operand : llvm::enumerate(op.getOperation()->getOperands())) {
       auto definingOp =
           dyn_cast_or_null<GenericOp>(operand.value().getDefiningOp());
       if (!definingOp || !definingOp.hasTensorSemantics())
         continue;
       auto fusedOp =
           fuseTensorOps(rewriter, cast<LinalgOp>(definingOp.getOperation()),
                         cast<LinalgOp>(op.getOperation()), operand.index());
       if (!fusedOp)
         continue;
       rewriter.replaceOp(op, fusedOp.getValue().getOperation()->getResults());
-      return matchSuccess();
+      return success();
     }
-    return matchFailure();
+    return failure();
   }
 };
 
 /// Pass that fuses generic ops on tensors. Used only for testing.
 struct FusionOfTensorOpsPass : public OperationPass<FusionOfTensorOpsPass> {
   void runOnOperation() override {
     OwningRewritePatternList patterns;
     Operation *op = getOperation();

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

 }
 
 template <typename LoopType, typename IndexedValueType, typename ConcreteOp>
 class LinalgRewritePattern : public RewritePattern {
 public:
   explicit LinalgRewritePattern(MLIRContext *context)
       : RewritePattern(ConcreteOp::getOperationName(), 1, context) {}
 
-  PatternMatchResult matchAndRewrite(Operation *op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(Operation *op,
     using Impl = LinalgOpToLoopsImpl<LoopType, IndexedValueType, ConcreteOp>;
     if (failed(Impl::doit(op, rewriter)))
-      return matchFailure();
+      return failure();
     rewriter.eraseOp(op);
-    return matchSuccess();
+    return success();
   }
 };
 
 // Helper classes for type list expansion.
 template <typename LoopType, typename IndexedValueType, typename... LinalgOps>
 class RewritePatternList;
 
 template <typename LoopType, typename IndexedValueType>

 //   3. One operands + returns a single AffineSymbolExpr
 //
 // In the first case, the AffineApplyOp is replaced by a new constant. In the
 // other cases, it is replaced by its unique operand.
 struct FoldAffineOp : public RewritePattern {
   FoldAffineOp(MLIRContext *context)
       : RewritePattern(AffineApplyOp::getOperationName(), 0, context) {}
 
-  PatternMatchResult matchAndRewrite(Operation *op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(Operation *op,
     AffineApplyOp affineApplyOp = cast<AffineApplyOp>(op);
     auto map = affineApplyOp.getAffineMap();
     if (map.getNumResults() != 1 || map.getNumInputs() > 1)
-      return matchFailure();
+      return failure();
 
     AffineExpr expr = map.getResult(0);
     if (map.getNumInputs() == 0) {
       if (auto val = expr.dyn_cast<AffineConstantExpr>()) {
         rewriter.replaceOpWithNewOp<ConstantIndexOp>(op, val.getValue());
-        return matchSuccess();
+        return success();
       }
-      return matchFailure();
+      return failure();
     }
     if (expr.dyn_cast<AffineDimExpr>() || expr.dyn_cast<AffineSymbolExpr>()) {
       rewriter.replaceOp(op, op->getOperand(0));
-      return matchSuccess();
+      return success();
     }
-    return matchFailure();
+    return failure();
   }
 };
 } // namespace
 
 template <typename LoopType, typename IndexedValueType>
 void LowerLinalgToLoopsPass<LoopType, IndexedValueType>::runOnFunction() {
   auto *context = &this->getContext();
   OwningRewritePatternList patterns;

mlir/lib/Dialect/QuantOps/Transforms/ConvertConst.cpp

 class ConvertConstPass : public FunctionPass<ConvertConstPass> {
 public:
   void runOnFunction() override;
 };
 
 struct QuantizedConstRewrite : public OpRewritePattern<QuantizeCastOp> {
   using OpRewritePattern<QuantizeCastOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(QuantizeCastOp qbarrier,
-                                     PatternRewriter &rewriter) const override;
+  LogicalResult matchAndRewrite(QuantizeCastOp qbarrier,
 };
 
 } // end anonymous namespace
 
 /// Matches a [constant] -> [qbarrier] where the qbarrier results type is
 /// quantized and the operand type is quantizable.
 
-PatternMatchResult
+LogicalResult
 QuantizedConstRewrite::matchAndRewrite(QuantizeCastOp qbarrier,
                                        PatternRewriter &rewriter) const {
   Attribute value;
 
   // Is the operand a constant?
   if (!matchPattern(qbarrier.arg(), m_Constant(&value))) {
-    return matchFailure();
+    return failure();
   }
 
   // Does the qbarrier convert to a quantized type. This will not be true
   // if a quantized type has not yet been chosen or if the cast to an equivalent
   // storage type is not supported.
   Type qbarrierResultType = qbarrier.getResult().getType();
   QuantizedType quantizedElementType =
       QuantizedType::getQuantizedElementType(qbarrierResultType);
   if (!quantizedElementType) {
-    return matchFailure();
+    return failure();
   }
   if (!QuantizedType::castToStorageType(qbarrierResultType)) {
-    return matchFailure();
+    return failure();
   }
 
   // Is the operand type compatible with the expressed type of the quantized
   // type? This will not be true if the qbarrier is superfluous (converts
   // from and to a quantized type).
   if (!quantizedElementType.isCompatibleExpressedType(
           qbarrier.arg().getType())) {
-    return matchFailure();
+    return failure();
   }
 
   // Is the constant value a type expressed in a way that we support?
   if (!value.isa<FloatAttr>() && !value.isa<DenseElementsAttr>() &&
       !value.isa<SparseElementsAttr>()) {
-    return matchFailure();
+    return failure();
   }
 
   Type newConstValueType;
   auto newConstValue =
       quantizeAttr(value, quantizedElementType, newConstValueType);
   if (!newConstValue) {
-    return matchFailure();
+    return failure();
   }
 
   // When creating the new const op, use a fused location that combines the
   // original const and the qbarrier that led to the quantization.
   auto fusedLoc = FusedLoc::get(
       {qbarrier.arg().getDefiningOp()->getLoc(), qbarrier.getLoc()},
       rewriter.getContext());
   auto newConstOp =
       rewriter.create<ConstantOp>(fusedLoc, newConstValueType, newConstValue);
   rewriter.replaceOpWithNewOp<StorageCastOp>(qbarrier, qbarrier.getType(),
                                              newConstOp);
-  return matchSuccess();
+  return success();
 }
 
 void ConvertConstPass::runOnFunction() {
   OwningRewritePatternList patterns;
   auto func = getFunction();
   auto *context = &getContext();
   patterns.insert<QuantizedConstRewrite>(context);
   applyPatternsGreedily(func, patterns);

mlir/lib/Dialect/QuantOps/Transforms/ConvertSimQuant.cpp

 template <typename ConcreteRewriteClass, typename FakeQuantOp>
 class FakeQuantRewrite : public OpRewritePattern<FakeQuantOp> {
 public:
   using OpRewritePattern<FakeQuantOp>::OpRewritePattern;
 
   FakeQuantRewrite(MLIRContext *ctx, bool *hadFailure)
       : OpRewritePattern<FakeQuantOp>(ctx), hadFailure(hadFailure) {}
 
-  PatternMatchResult matchAndRewrite(FakeQuantOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(FakeQuantOp op,
     // TODO: If this pattern comes up more frequently, consider adding core
     // support for failable rewrites.
     if (failableRewrite(op, rewriter)) {
       *hadFailure = true;
-      return Pattern::matchFailure();
+      return failure();
     }
 
-    return Pattern::matchSuccess();
+    return success();
   }
 
 private:
   bool *hadFailure;
 
   bool failableRewrite(FakeQuantOp op, PatternRewriter &rewriter) const {
     auto converter = ExpressedToQuantizedConverter::forInputType(op.getType());
     if (!converter) {

mlir/lib/Dialect/SPIRV/SPIRVCanonicalization.cpp

 namespace {
 
 /// Combines chained `spirv::AccessChainOp` operations into one
 /// `spirv::AccessChainOp` operation.
 struct CombineChainedAccessChain
     : public OpRewritePattern<spirv::AccessChainOp> {
   using OpRewritePattern<spirv::AccessChainOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(spirv::AccessChainOp accessChainOp,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(spirv::AccessChainOp accessChainOp,
     auto parentAccessChainOp = dyn_cast_or_null<spirv::AccessChainOp>(
         accessChainOp.base_ptr().getDefiningOp());
 
     if (!parentAccessChainOp) {
-      return matchFailure();
+      return failure();
     }
 
     // Combine indices.
     SmallVector<Value, 4> indices(parentAccessChainOp.indices());
     indices.append(accessChainOp.indices().begin(),
                    accessChainOp.indices().end());
 
     rewriter.replaceOpWithNewOp<spirv::AccessChainOp>(
         accessChainOp, parentAccessChainOp.base_ptr(), indices);
 
-    return matchSuccess();
+    return success();
   }
 };
 } // end anonymous namespace
 
 void spirv::AccessChainOp::getCanonicalizationPatterns(
     OwningRewritePatternList &results, MLIRContext *context) {
   results.insert<CombineChainedAccessChain>(context);
 }

 //                       +-------------+
 //                       | merge block |
 //                       +-------------+
 //
 struct ConvertSelectionOpToSelect
     : public OpRewritePattern<spirv::SelectionOp> {
   using OpRewritePattern<spirv::SelectionOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(spirv::SelectionOp selectionOp,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(spirv::SelectionOp selectionOp,
     auto *op = selectionOp.getOperation();
     auto &body = op->getRegion(0);
     // Verifier allows an empty region for `spv.selection`.
     if (body.empty()) {
-      return matchFailure();
+      return failure();
     }
 
     // Check that region consists of 4 blocks:
     // header block, `true` block, `false` block and merge block.
     if (std::distance(body.begin(), body.end()) != 4) {
-      return matchFailure();
+      return failure();
     }
 
     auto *headerBlock = selectionOp.getHeaderBlock();
     if (!onlyContainsBranchConditionalOp(headerBlock)) {
-      return matchFailure();
+      return failure();
     }
 
     auto brConditionalOp =
         cast<spirv::BranchConditionalOp>(headerBlock->front());
 
     auto *trueBlock = brConditionalOp.getSuccessor(0);
     auto *falseBlock = brConditionalOp.getSuccessor(1);
     auto *mergeBlock = selectionOp.getMergeBlock();
 
     if (failed(canCanonicalizeSelection(trueBlock, falseBlock, mergeBlock)))
-      return matchFailure();
+      return failure();
 
     auto trueValue = getSrcValue(trueBlock);
     auto falseValue = getSrcValue(falseBlock);
     auto ptrValue = getDstPtr(trueBlock);
     auto storeOpAttributes =
         cast<spirv::StoreOp>(trueBlock->front()).getOperation()->getAttrs();
 
     auto selectOp = rewriter.create<spirv::SelectOp>(
         selectionOp.getLoc(), trueValue.getType(), brConditionalOp.condition(),
         trueValue, falseValue);
     rewriter.create<spirv::StoreOp>(selectOp.getLoc(), ptrValue,
                                     selectOp.getResult(), storeOpAttributes);
 
     // `spv.selection` is not needed anymore.
     rewriter.eraseOp(op);
-    return matchSuccess();
+    return success();
   }
 
 private:
   // Checks that given blocks follow the following rules:
   // 1. Each conditional block consists of two operations, the first operation
   //    is a `spv.Store` and the last operation is a `spv.Branch`.
   // 2. Each `spv.Store` uses the same pointer and the same memory attributes.
   // 3. A control flow goes into the given merge block from the given
   //    conditional blocks.
-  PatternMatchResult canCanonicalizeSelection(Block *trueBlock,
-                                              Block *falseBlock,
-                                              Block *mergeBlock) const;
+  LogicalResult canCanonicalizeSelection(Block *trueBlock, Block *falseBlock,
 
   bool onlyContainsBranchConditionalOp(Block *block) const {
     return std::next(block->begin()) == block->end() &&
            isa<spirv::BranchConditionalOp>(block->front());
   }
 
   bool isSameAttrList(spirv::StoreOp lhs, spirv::StoreOp rhs) const {
     return lhs.getOperation()->getAttrList().getDictionary() ==
private:
 
   // Returns a destination value for the given block.
   Value getDstPtr(Block *block) const {
     auto storeOp = cast<spirv::StoreOp>(block->front());
     return storeOp.ptr();
   }
 };
 
-PatternMatchResult ConvertSelectionOpToSelect::canCanonicalizeSelection(
+LogicalResult ConvertSelectionOpToSelect::canCanonicalizeSelection(
     Block *trueBlock, Block *falseBlock, Block *mergeBlock) const {
   // Each block must consists of 2 operations.
   if ((std::distance(trueBlock->begin(), trueBlock->end()) != 2) ||
       (std::distance(falseBlock->begin(), falseBlock->end()) != 2)) {
-    return matchFailure();
+    return failure();
   }
 
   auto trueBrStoreOp = dyn_cast<spirv::StoreOp>(trueBlock->front());
   auto trueBrBranchOp =
       dyn_cast<spirv::BranchOp>(*std::next(trueBlock->begin()));
   auto falseBrStoreOp = dyn_cast<spirv::StoreOp>(falseBlock->front());
   auto falseBrBranchOp =
       dyn_cast<spirv::BranchOp>(*std::next(falseBlock->begin()));
 
   if (!trueBrStoreOp || !trueBrBranchOp || !falseBrStoreOp ||
       !falseBrBranchOp) {
-    return matchFailure();
+    return failure();
   }
 
   // Check that each `spv.Store` uses the same pointer, memory access
   // attributes and a valid type of the value.
   if ((trueBrStoreOp.ptr() != falseBrStoreOp.ptr()) ||
       !isSameAttrList(trueBrStoreOp, falseBrStoreOp) ||
       !isValidType(trueBrStoreOp.value().getType())) {
-    return matchFailure();
+    return failure();
   }
 
   if ((trueBrBranchOp.getOperation()->getSuccessor(0) != mergeBlock) ||
       (falseBrBranchOp.getOperation()->getSuccessor(0) != mergeBlock)) {
-    return matchFailure();
+    return failure();
   }
 
-  return matchSuccess();
+  return success();
 }
 } // end anonymous namespace
 
 void spirv::SelectionOp::getCanonicalizationPatterns(
     OwningRewritePatternList &results, MLIRContext *context) {
   results.insert<ConvertSelectionOpToSelect>(context);
 }

mlir/lib/Dialect/SPIRV/SPIRVLowering.cpp

 
 namespace {
 /// A pattern for rewriting function signature to convert arguments of functions
 /// to be of valid SPIR-V types.
 class FuncOpConversion final : public SPIRVOpLowering<FuncOp> {
 public:
   using SPIRVOpLowering<FuncOp>::SPIRVOpLowering;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(FuncOp funcOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override;
 };
 } // namespace
 
-PatternMatchResult
+LogicalResult
 FuncOpConversion::matchAndRewrite(FuncOp funcOp, ArrayRef<Value> operands,
                                   ConversionPatternRewriter &rewriter) const {
   auto fnType = funcOp.getType();
   // TODO(antiagainst): support converting functions with one result.
   if (fnType.getNumResults())
-    return matchFailure();
+    return failure();
 
   TypeConverter::SignatureConversion signatureConverter(fnType.getNumInputs());
   for (auto argType : enumerate(funcOp.getType().getInputs())) {
     auto convertedType = typeConverter.convertType(argType.value());
     if (!convertedType)
-      return matchFailure();
+      return failure();
     signatureConverter.addInputs(argType.index(), convertedType);
   }
 
   // Create the converted spv.func op.
   auto newFuncOp = rewriter.create<spirv::FuncOp>(
       funcOp.getLoc(), funcOp.getName(),
       rewriter.getFunctionType(signatureConverter.getConvertedTypes(),
                                llvm::None));
 
   // Copy over all attributes other than the function name and type.
   for (const auto &namedAttr : funcOp.getAttrs()) {
     if (!namedAttr.first.is(impl::getTypeAttrName()) &&
         !namedAttr.first.is(SymbolTable::getSymbolAttrName()))
       newFuncOp.setAttr(namedAttr.first, namedAttr.second);
   }
 
   rewriter.inlineRegionBefore(funcOp.getBody(), newFuncOp.getBody(),
                               newFuncOp.end());
   rewriter.applySignatureConversion(&newFuncOp.getBody(), signatureConverter);
   rewriter.eraseOp(funcOp);
-  return matchSuccess();
+  return success();
 }
 
 void mlir::populateBuiltinFuncToSPIRVPatterns(
     MLIRContext *context, SPIRVTypeConverter &typeConverter,
     OwningRewritePatternList &patterns) {
   patterns.insert<FuncOpConversion>(context, typeConverter);
 }
 

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

 using namespace mlir;
 
 namespace {
 class SPIRVGlobalVariableOpLayoutInfoDecoration
     : public OpRewritePattern<spirv::GlobalVariableOp> {
 public:
   using OpRewritePattern<spirv::GlobalVariableOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(spirv::GlobalVariableOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(spirv::GlobalVariableOp op,
     spirv::StructType::LayoutInfo structSize = 0;
     VulkanLayoutUtils::Size structAlignment = 1;
     SmallVector<NamedAttribute, 4> globalVarAttrs;
 
     auto ptrType = op.type().cast<spirv::PointerType>();
     auto structType = VulkanLayoutUtils::decorateType(
         ptrType.getPointeeType().cast<spirv::StructType>(), structSize,
         structAlignment);
     auto decoratedType =
         spirv::PointerType::get(structType, ptrType.getStorageClass());
 
     // Save all named attributes except "type" attribute.
     for (const auto &attr : op.getAttrs()) {
       if (attr.first == "type") {
         continue;
       }
       globalVarAttrs.push_back(attr);
     }
 
     rewriter.replaceOpWithNewOp<spirv::GlobalVariableOp>(
         op, TypeAttr::get(decoratedType), globalVarAttrs);
-    return matchSuccess();
+    return success();
   }
 };
 
 class SPIRVAddressOfOpLayoutInfoDecoration
     : public OpRewritePattern<spirv::AddressOfOp> {
 public:
   using OpRewritePattern<spirv::AddressOfOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(spirv::AddressOfOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(spirv::AddressOfOp op,
     auto spirvModule = op.getParentOfType<spirv::ModuleOp>();
     auto varName = op.variable();
     auto varOp = spirvModule.lookupSymbol<spirv::GlobalVariableOp>(varName);
 
     rewriter.replaceOpWithNewOp<spirv::AddressOfOp>(
         op, varOp.type(), rewriter.getSymbolRefAttr(varName));
-    return matchSuccess();
+    return success();
   }
 };
 } // namespace
 
 static void populateSPIRVLayoutInfoPatterns(OwningRewritePatternList &patterns,
                                             MLIRContext *ctx) {
   patterns.insert<SPIRVGlobalVariableOpLayoutInfoDecoration,
                   SPIRVAddressOfOpLayoutInfoDecoration>(ctx);

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

 ///
 /// Specifically, this pattern creates global variables according to interface
 /// variable ABI attributes attached to function arguments and converts all
 /// function argument uses to those global variables. This is necessary because
 /// Vulkan requires all shader entry points to be of void(void) type.
 class ProcessInterfaceVarABI final : public SPIRVOpLowering<spirv::FuncOp> {
 public:
   using SPIRVOpLowering<spirv::FuncOp>::SPIRVOpLowering;
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(spirv::FuncOp funcOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
 /// Pass to implement the ABI information specified as attributes.
 class LowerABIAttributesPass final
     : public OperationPass<LowerABIAttributesPass, spirv::ModuleOp> {
 private:
   void runOnOperation() override;
 };
 } // namespace
 
-PatternMatchResult ProcessInterfaceVarABI::matchAndRewrite(
+LogicalResult ProcessInterfaceVarABI::matchAndRewrite(
     spirv::FuncOp funcOp, ArrayRef<Value> operands,
     ConversionPatternRewriter &rewriter) const {
   if (!funcOp.getAttrOfType<spirv::EntryPointABIAttr>(
           spirv::getEntryPointABIAttrName())) {
     // TODO(ravishankarm) : Non-entry point functions are not handled.
-    return matchFailure();
+    return failure();
   }
   TypeConverter::SignatureConversion signatureConverter(
       funcOp.getType().getNumInputs());
 
   auto attrName = spirv::getInterfaceVarABIAttrName();
   for (auto argType : llvm::enumerate(funcOp.getType().getInputs())) {
     auto abiInfo = funcOp.getArgAttrOfType<spirv::InterfaceVarABIAttr>(
         argType.index(), attrName);
     if (!abiInfo) {
       // TODO(ravishankarm) : For non-entry point functions, it should be legal
       // to pass around scalar/vector values and return a scalar/vector. For now
       // non-entry point functions are not handled in this ABI lowering and will
       // produce an error.
-      return matchFailure();
+      return failure();
     }
     auto var =
         createGlobalVariableForArg(funcOp, rewriter, argType.index(), abiInfo);
     if (!var) {
-      return matchFailure();
+      return failure();
     }
 
     OpBuilder::InsertionGuard funcInsertionGuard(rewriter);
     rewriter.setInsertionPointToStart(&funcOp.front());
     // Insert spirv::AddressOf and spirv::AccessChain operations.
     Value replacement =
         rewriter.create<spirv::AddressOfOp>(funcOp.getLoc(), var);
     // Check if the arg is a scalar or vector type. In that case, the value
LogicalResult ProcessInterfaceVarABI::matchAndRewrite(
   }
 
   // Creates a new function with the update signature.
   rewriter.updateRootInPlace(funcOp, [&] {
     funcOp.setType(rewriter.getFunctionType(
         signatureConverter.getConvertedTypes(), llvm::None));
     rewriter.applySignatureConversion(&funcOp.getBody(), signatureConverter);
   });
-  return matchSuccess();
+  return success();
 }
 
 void LowerABIAttributesPass::runOnOperation() {
   // Uses the signature conversion methodology of the dialect conversion
   // framework to implement the conversion.
   spirv::ModuleOp module = getOperation();
   MLIRContext *context = &getContext();
 

mlir/lib/Dialect/StandardOps/IR/Ops.cpp

static LogicalResult verify(AllocOp op) {
   return success();
 }
 
 namespace {
 /// Fold constant dimensions into an alloc operation.
 struct SimplifyAllocConst : public OpRewritePattern<AllocOp> {
   using OpRewritePattern<AllocOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(AllocOp alloc,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(AllocOp alloc,
     // Check to see if any dimensions operands are constants.  If so, we can
     // substitute and drop them.
     if (llvm::none_of(alloc.getOperands(), [](Value operand) {
           return matchPattern(operand, m_ConstantIndex());
         }))
-      return matchFailure();
+      return failure();
 
     auto memrefType = alloc.getType();
 
     // Ok, we have one or more constant operands.  Collect the non-constant ones
     // and keep track of the resultant memref type to build.
     SmallVector<int64_t, 4> newShapeConstants;
     newShapeConstants.reserve(memrefType.getRank());
     SmallVector<Value, 4> newOperands;
LogicalResult matchAndRewrite(AllocOp alloc,
     // Create and insert the alloc op for the new memref.
     auto newAlloc = rewriter.create<AllocOp>(alloc.getLoc(), newMemRefType,
                                              newOperands, IntegerAttr());
     // Insert a cast so we have the same type as the old alloc.
     auto resultCast = rewriter.create<MemRefCastOp>(alloc.getLoc(), newAlloc,
                                                     alloc.getType());
 
     rewriter.replaceOp(alloc, {resultCast});
-    return matchSuccess();
+    return success();
   }
 };
 
 /// Fold alloc operations with no uses. Alloc has side effects on the heap,
 /// but can still be deleted if it has zero uses.
 struct SimplifyDeadAlloc : public OpRewritePattern<AllocOp> {
   using OpRewritePattern<AllocOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(AllocOp alloc,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(AllocOp alloc,
     if (alloc.use_empty()) {
       rewriter.eraseOp(alloc);
-      return matchSuccess();
+      return success();
     }
-    return matchFailure();
+    return failure();
   }
 };
 } // end anonymous namespace.
 
 void AllocOp::getCanonicalizationPatterns(OwningRewritePatternList &results,
                                           MLIRContext *context) {
   results.insert<SimplifyAllocConst, SimplifyDeadAlloc>(context);
 }

 //===----------------------------------------------------------------------===//
 
 namespace {
 /// Simplify a branch to a block that has a single predecessor. This effectively
 /// merges the two blocks.
 struct SimplifyBrToBlockWithSinglePred : public OpRewritePattern<BranchOp> {
   using OpRewritePattern<BranchOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(BranchOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(BranchOp op,
     // Check that the successor block has a single predecessor.
     Block *succ = op.getDest();
     Block *opParent = op.getOperation()->getBlock();
     if (succ == opParent || !has_single_element(succ->getPredecessors()))
-      return matchFailure();
+      return failure();
 
     // Merge the successor into the current block and erase the branch.
     rewriter.mergeBlocks(succ, opParent, op.getOperands());
     rewriter.eraseOp(op);
-    return matchSuccess();
+    return success();
   }
 };
 } // end anonymous namespace.
 
 Block *BranchOp::getDest() { return getSuccessor(); }
 
 void BranchOp::setDest(Block *block) { return setSuccessor(block); }
 

 // CallIndirectOp
 //===----------------------------------------------------------------------===//
 namespace {
 /// Fold indirect calls that have a constant function as the callee operand.
 struct SimplifyIndirectCallWithKnownCallee
     : public OpRewritePattern<CallIndirectOp> {
   using OpRewritePattern<CallIndirectOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(CallIndirectOp indirectCall,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(CallIndirectOp indirectCall,
     // Check that the callee is a constant callee.
     SymbolRefAttr calledFn;
     if (!matchPattern(indirectCall.getCallee(), m_Constant(&calledFn)))
-      return matchFailure();
+      return failure();
 
     // Replace with a direct call.
     rewriter.replaceOpWithNewOp<CallOp>(indirectCall, calledFn,
                                         indirectCall.getResultTypes(),
                                         indirectCall.getArgOperands());
-    return matchSuccess();
+    return success();
   }
 };
 } // end anonymous namespace.
 
 void CallIndirectOp::getCanonicalizationPatterns(
     OwningRewritePatternList &results, MLIRContext *context) {
   results.insert<SimplifyIndirectCallWithKnownCallee>(context);
 }

 
 namespace {
 /// cond_br true, ^bb1, ^bb2 -> br ^bb1
 /// cond_br false, ^bb1, ^bb2 -> br ^bb2
 ///
 struct SimplifyConstCondBranchPred : public OpRewritePattern<CondBranchOp> {
   using OpRewritePattern<CondBranchOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(CondBranchOp condbr,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(CondBranchOp condbr,
     if (matchPattern(condbr.getCondition(), m_NonZero())) {
       // True branch taken.
       rewriter.replaceOpWithNewOp<BranchOp>(condbr, condbr.getTrueDest(),
                                             condbr.getTrueOperands());
-      return matchSuccess();
+      return success();
     } else if (matchPattern(condbr.getCondition(), m_Zero())) {
       // False branch taken.
       rewriter.replaceOpWithNewOp<BranchOp>(condbr, condbr.getFalseDest(),
                                             condbr.getFalseOperands());
-      return matchSuccess();
+      return success();
     }
-    return matchFailure();
+    return failure();
   }
 };
 } // end anonymous namespace.
 
 void CondBranchOp::getCanonicalizationPatterns(
     OwningRewritePatternList &results, MLIRContext *context) {
   results.insert<SimplifyConstCondBranchPred>(context);
 }

 // DeallocOp
 //===----------------------------------------------------------------------===//
 namespace {
 /// Fold Dealloc operations that are deallocating an AllocOp that is only used
 /// by other Dealloc operations.
 struct SimplifyDeadDealloc : public OpRewritePattern<DeallocOp> {
   using OpRewritePattern<DeallocOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(DeallocOp dealloc,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(DeallocOp dealloc,
     // Check that the memref operand's defining operation is an AllocOp.
     Value memref = dealloc.memref();
     if (!isa_and_nonnull<AllocOp>(memref.getDefiningOp()))
-      return matchFailure();
+      return failure();
 
     // Check that all of the uses of the AllocOp are other DeallocOps.
     for (auto *user : memref.getUsers())
       if (!isa<DeallocOp>(user))
-        return matchFailure();
+        return failure();
 
     // Erase the dealloc operation.
     rewriter.eraseOp(dealloc);
-    return matchSuccess();
+    return success();
   }
 };
 } // end anonymous namespace.
 
 static LogicalResult verify(DeallocOp op) {
   if (!op.memref().getType().isa<MemRefType>())
     return op.emitOpError("operand must be a memref");
   return success();

 
 namespace {
 
 /// Pattern to rewrite a subview op with constant size arguments.
 class SubViewOpShapeFolder final : public OpRewritePattern<SubViewOp> {
 public:
   using OpRewritePattern<SubViewOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(SubViewOp subViewOp,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(SubViewOp subViewOp,
     MemRefType subViewType = subViewOp.getType();
     // Follow all or nothing approach for shapes for now. If all the operands
     // for sizes are constants then fold it into the type of the result memref.
     if (subViewType.hasStaticShape() ||
         llvm::any_of(subViewOp.sizes(), [](Value operand) {
           return !matchPattern(operand, m_ConstantIndex());
         })) {
-      return matchFailure();
+      return failure();
     }
     SmallVector<int64_t, 4> staticShape(subViewOp.getNumSizes());
     for (auto size : llvm::enumerate(subViewOp.sizes())) {
       auto defOp = size.value().getDefiningOp();
       assert(defOp);
       staticShape[size.index()] = cast<ConstantIndexOp>(defOp).getValue();
     }
     MemRefType newMemRefType =
         MemRefType::Builder(subViewType).setShape(staticShape);
     auto newSubViewOp = rewriter.create<SubViewOp>(
         subViewOp.getLoc(), subViewOp.source(), subViewOp.offsets(),
         ArrayRef<Value>(), subViewOp.strides(), newMemRefType);
     // Insert a memref_cast for compatibility of the uses of the op.
     rewriter.replaceOpWithNewOp<MemRefCastOp>(subViewOp, newSubViewOp,
                                               subViewOp.getType());
-    return matchSuccess();
+    return success();
   }
 };
 
 // Pattern to rewrite a subview op with constant stride arguments.
 class SubViewOpStrideFolder final : public OpRewritePattern<SubViewOp> {
 public:
   using OpRewritePattern<SubViewOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(SubViewOp subViewOp,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(SubViewOp subViewOp,
     if (subViewOp.getNumStrides() == 0) {
-      return matchFailure();
+      return failure();
     }
     // Follow all or nothing approach for strides for now. If all the operands
     // for strides are constants then fold it into the strides of the result
     // memref.
     int64_t baseOffset, resultOffset;
     SmallVector<int64_t, 4> baseStrides, resultStrides;
     MemRefType subViewType = subViewOp.getType();
     if (failed(getStridesAndOffset(subViewOp.getBaseMemRefType(), baseStrides,
                                    baseOffset)) ||
         failed(getStridesAndOffset(subViewType, resultStrides, resultOffset)) ||
         llvm::is_contained(baseStrides,
                            MemRefType::getDynamicStrideOrOffset()) ||
         llvm::any_of(subViewOp.strides(), [](Value stride) {
           return !matchPattern(stride, m_ConstantIndex());
         })) {
-      return matchFailure();
+      return failure();
     }
 
     SmallVector<int64_t, 4> staticStrides(subViewOp.getNumStrides());
     for (auto stride : llvm::enumerate(subViewOp.strides())) {
       auto defOp = stride.value().getDefiningOp();
       assert(defOp);
       assert(baseStrides[stride.index()] > 0);
       staticStrides[stride.index()] =
           cast<ConstantIndexOp>(defOp).getValue() * baseStrides[stride.index()];
     }
     AffineMap layoutMap = makeStridedLinearLayoutMap(
         staticStrides, resultOffset, rewriter.getContext());
     MemRefType newMemRefType =
         MemRefType::Builder(subViewType).setAffineMaps(layoutMap);
     auto newSubViewOp = rewriter.create<SubViewOp>(
         subViewOp.getLoc(), subViewOp.source(), subViewOp.offsets(),
         subViewOp.sizes(), ArrayRef<Value>(), newMemRefType);
     // Insert a memref_cast for compatibility of the uses of the op.
     rewriter.replaceOpWithNewOp<MemRefCastOp>(subViewOp, newSubViewOp,
                                               subViewOp.getType());
-    return matchSuccess();
+    return success();
   }
 };
 
 // Pattern to rewrite a subview op with constant offset arguments.
 class SubViewOpOffsetFolder final : public OpRewritePattern<SubViewOp> {
 public:
   using OpRewritePattern<SubViewOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(SubViewOp subViewOp,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(SubViewOp subViewOp,
     if (subViewOp.getNumOffsets() == 0) {
-      return matchFailure();
+      return failure();
     }
     // Follow all or nothing approach for offsets for now. If all the operands
     // for offsets are constants then fold it into the offset of the result
     // memref.
     int64_t baseOffset, resultOffset;
     SmallVector<int64_t, 4> baseStrides, resultStrides;
     MemRefType subViewType = subViewOp.getType();
     if (failed(getStridesAndOffset(subViewOp.getBaseMemRefType(), baseStrides,
                                    baseOffset)) ||
         failed(getStridesAndOffset(subViewType, resultStrides, resultOffset)) ||
         llvm::is_contained(baseStrides,
                            MemRefType::getDynamicStrideOrOffset()) ||
         baseOffset == MemRefType::getDynamicStrideOrOffset() ||
         llvm::any_of(subViewOp.offsets(), [](Value stride) {
           return !matchPattern(stride, m_ConstantIndex());
         })) {
-      return matchFailure();
+      return failure();
     }
 
     auto staticOffset = baseOffset;
     for (auto offset : llvm::enumerate(subViewOp.offsets())) {
       auto defOp = offset.value().getDefiningOp();
       assert(defOp);
       assert(baseStrides[offset.index()] > 0);
       staticOffset +=
           cast<ConstantIndexOp>(defOp).getValue() * baseStrides[offset.index()];
     }
 
     AffineMap layoutMap = makeStridedLinearLayoutMap(
         resultStrides, staticOffset, rewriter.getContext());
     MemRefType newMemRefType =
         MemRefType::Builder(subViewType).setAffineMaps(layoutMap);
     auto newSubViewOp = rewriter.create<SubViewOp>(
         subViewOp.getLoc(), subViewOp.source(), ArrayRef<Value>(),
         subViewOp.sizes(), subViewOp.strides(), newMemRefType);
     // Insert a memref_cast for compatibility of the uses of the op.
     rewriter.replaceOpWithNewOp<MemRefCastOp>(subViewOp, newSubViewOp,
                                               subViewOp.getType());
-    return matchSuccess();
+    return success();
   }
 };
 
 } // end anonymous namespace
 
 void SubViewOp::getCanonicalizationPatterns(OwningRewritePatternList &results,
                                             MLIRContext *context) {
   results.insert<SubViewOpShapeFolder, SubViewOpStrideFolder,
static LogicalResult verify(ViewOp op) {
   return success();
 }
 
 namespace {
 
 struct ViewOpShapeFolder : public OpRewritePattern<ViewOp> {
   using OpRewritePattern<ViewOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(ViewOp viewOp,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(ViewOp viewOp,
     // Return if none of the operands are constants.
     if (llvm::none_of(viewOp.getOperands(), [](Value operand) {
           return matchPattern(operand, m_ConstantIndex());
         }))
-      return matchFailure();
+      return failure();
 
     // Get result memref type.
     auto memrefType = viewOp.getType();
     if (memrefType.getAffineMaps().size() > 1)
-      return matchFailure();
+      return failure();
     auto map = memrefType.getAffineMaps().empty()
                    ? AffineMap::getMultiDimIdentityMap(memrefType.getRank(),
                                                        rewriter.getContext())
                    : memrefType.getAffineMaps()[0];
 
     // Get offset from old memref view type 'memRefType'.
     int64_t oldOffset;
     SmallVector<int64_t, 4> oldStrides;
     if (failed(getStridesAndOffset(memrefType, oldStrides, oldOffset)))
-      return matchFailure();
+      return failure();
 
     SmallVector<Value, 4> newOperands;
 
     // Fold dynamic offset operand if it is produced by a constant.
     auto dynamicOffset = viewOp.getDynamicOffset();
     int64_t newOffset = oldOffset;
     unsigned dynamicOffsetOperandCount = 0;
     if (dynamicOffset != nullptr) {
assert(static_cast<int64_t>(newOperands.size()) ==
            dynamicOffsetOperandCount + newMemRefType.getNumDynamicDims());
 
     // Create new ViewOp.
     auto newViewOp = rewriter.create<ViewOp>(viewOp.getLoc(), newMemRefType,
                                              viewOp.getOperand(0), newOperands);
     // Insert a cast so we have the same type as the old memref type.
     rewriter.replaceOpWithNewOp<MemRefCastOp>(viewOp, newViewOp,
                                               viewOp.getType());
-    return matchSuccess();
+    return success();
   }
 };
 
 struct ViewOpMemrefCastFolder : public OpRewritePattern<ViewOp> {
   using OpRewritePattern<ViewOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(ViewOp viewOp,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(ViewOp viewOp,
     Value memrefOperand = viewOp.getOperand(0);
     MemRefCastOp memrefCastOp =
         dyn_cast_or_null<MemRefCastOp>(memrefOperand.getDefiningOp());
     if (!memrefCastOp)
-      return matchFailure();
+      return failure();
     Value allocOperand = memrefCastOp.getOperand();
     AllocOp allocOp = dyn_cast_or_null<AllocOp>(allocOperand.getDefiningOp());
     if (!allocOp)
-      return matchFailure();
+      return failure();
     rewriter.replaceOpWithNewOp<ViewOp>(viewOp, viewOp.getType(), allocOperand,
                                         viewOp.operands());
-    return matchSuccess();
+    return success();
   }
 };
 
 } // end anonymous namespace
 
 void ViewOp::getCanonicalizationPatterns(OwningRewritePatternList &results,
                                          MLIRContext *context) {
   results.insert<ViewOpShapeFolder, ViewOpMemrefCastFolder>(context);

mlir/lib/Dialect/VectorOps/VectorOps.cpp

 namespace {
 
 // Pattern to rewrite a StridedSliceOp(ConstantMaskOp) -> ConstantMaskOp.
 class StridedSliceConstantMaskFolder final
     : public OpRewritePattern<StridedSliceOp> {
 public:
   using OpRewritePattern<StridedSliceOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(StridedSliceOp stridedSliceOp,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(StridedSliceOp stridedSliceOp,
     // Return if 'stridedSliceOp' operand is not defined by a ConstantMaskOp.
     auto defOp = stridedSliceOp.vector().getDefiningOp();
     auto constantMaskOp = dyn_cast_or_null<ConstantMaskOp>(defOp);
     if (!constantMaskOp)
-      return matchFailure();
+      return failure();
     // Return if 'stridedSliceOp' has non-unit strides.
     if (llvm::any_of(stridedSliceOp.strides(), [](Attribute attr) {
           return attr.cast<IntegerAttr>().getInt() != 1;
         }))
-      return matchFailure();
+      return failure();
     // Gather constant mask dimension sizes.
     SmallVector<int64_t, 4> maskDimSizes;
     populateFromInt64AttrArray(constantMaskOp.mask_dim_sizes(), maskDimSizes);
     // Gather strided slice offsets and sizes.
     SmallVector<int64_t, 4> sliceOffsets;
     populateFromInt64AttrArray(stridedSliceOp.offsets(), sliceOffsets);
     SmallVector<int64_t, 4> sliceSizes;
     populateFromInt64AttrArray(stridedSliceOp.sizes(), sliceSizes);
LogicalResult matchAndRewrite(StridedSliceOp stridedSliceOp,
     // region is a conjunction of mask dim intervals).
     if (llvm::any_of(sliceMaskDimSizes, [](int64_t sz) { return sz == 0; }))
       sliceMaskDimSizes.assign(maskDimSizes.size(), 0);
 
     // Replace 'stridedSliceOp' with ConstantMaskOp with sliced mask region.
     rewriter.replaceOpWithNewOp<ConstantMaskOp>(
         stridedSliceOp, stridedSliceOp.getResult().getType(),
         vector::getVectorSubscriptAttr(rewriter, sliceMaskDimSizes));
-    return matchSuccess();
+    return success();
   }
 };
 
 } // end anonymous namespace
 
 void StridedSliceOp::getCanonicalizationPatterns(
     OwningRewritePatternList &results, MLIRContext *context) {
   // Pattern to rewrite a StridedSliceOp(ConstantMaskOp) -> ConstantMaskOp.

 
 namespace {
 
 // Pattern to rewrite a CreateMaskOp with a ConstantMaskOp.
 class CreateMaskFolder final : public OpRewritePattern<CreateMaskOp> {
 public:
   using OpRewritePattern<CreateMaskOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(CreateMaskOp createMaskOp,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(CreateMaskOp createMaskOp,
     // Return if any of 'createMaskOp' operands are not defined by a constant.
     auto is_not_def_by_constant = [](Value operand) {
       return !isa_and_nonnull<ConstantIndexOp>(operand.getDefiningOp());
     };
     if (llvm::any_of(createMaskOp.operands(), is_not_def_by_constant))
-      return matchFailure();
+      return failure();
     // Gather constant mask dimension sizes.
     SmallVector<int64_t, 4> maskDimSizes;
     for (auto operand : createMaskOp.operands()) {
       auto defOp = operand.getDefiningOp();
       maskDimSizes.push_back(cast<ConstantIndexOp>(defOp).getValue());
     }
     // Replace 'createMaskOp' with ConstantMaskOp.
     rewriter.replaceOpWithNewOp<ConstantMaskOp>(
         createMaskOp, createMaskOp.getResult().getType(),
         vector::getVectorSubscriptAttr(rewriter, maskDimSizes));
-    return matchSuccess();
+    return success();
   }
 };
 
 } // end anonymous namespace
 
 void CreateMaskOp::getCanonicalizationPatterns(
     OwningRewritePatternList &results, MLIRContext *context) {
   results.insert<CreateMaskFolder>(context);

mlir/lib/Dialect/VectorOps/VectorTransforms.cpp

 
 namespace {
 
 // Splits vector TransferReadOp into smaller TransferReadOps based on slicing
 // scheme of its unique ExtractSlicesOp user.
 struct SplitTransferReadOp : public OpRewritePattern<vector::TransferReadOp> {
   using OpRewritePattern<vector::TransferReadOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(vector::TransferReadOp xferReadOp,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(vector::TransferReadOp xferReadOp,
     // TODO(andydavis, ntv) Support splitting TransferReadOp with non-identity
     // permutation maps. Repurpose code from MaterializeVectors transformation.
     if (!isIdentitySuffix(xferReadOp.permutation_map()))
-      return matchFailure();
+      return failure();
     // Return unless the unique 'xferReadOp' user is an ExtractSlicesOp.
     Value xferReadResult = xferReadOp.getResult();
     auto extractSlicesOp =
         dyn_cast<vector::ExtractSlicesOp>(*xferReadResult.getUsers().begin());
     if (!xferReadResult.hasOneUse() || !extractSlicesOp)
-      return matchFailure();
+      return failure();
 
     // Get 'sizes' and 'strides' parameters from ExtractSlicesOp user.
     auto sourceVectorType = extractSlicesOp.getSourceVectorType();
     auto resultTupleType = extractSlicesOp.getResultTupleType();
     SmallVector<int64_t, 4> sizes;
     extractSlicesOp.getSizes(sizes);
     SmallVector<int64_t, 4> strides;
     extractSlicesOp.getStrides(strides);
LogicalResult matchAndRewrite(vector::TransferReadOp xferReadOp,
 
     // Create tuple of splice xfer read operations.
     Value tupleOp = rewriter.create<vector::TupleOp>(loc, resultTupleType,
                                                      vectorTupleValues);
     // Replace 'xferReadOp' with result 'insertSlicesResult'.
     rewriter.replaceOpWithNewOp<vector::InsertSlicesOp>(
         xferReadOp, sourceVectorType, tupleOp, extractSlicesOp.sizes(),
         extractSlicesOp.strides());
-    return matchSuccess();
+    return success();
   }
 };
 
 // Splits vector TransferWriteOp into smaller TransferWriteOps for each source.
 struct SplitTransferWriteOp : public OpRewritePattern<vector::TransferWriteOp> {
   using OpRewritePattern<vector::TransferWriteOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(vector::TransferWriteOp xferWriteOp,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(vector::TransferWriteOp xferWriteOp,
     // TODO(andydavis, ntv) Support splitting TransferWriteOp with non-identity
     // permutation maps. Repurpose code from MaterializeVectors transformation.
     if (!isIdentitySuffix(xferWriteOp.permutation_map()))
-      return matchFailure();
+      return failure();
     // Return unless the 'xferWriteOp' 'vector' operand is an 'InsertSlicesOp'.
     auto *vectorDefOp = xferWriteOp.vector().getDefiningOp();
     auto insertSlicesOp = dyn_cast_or_null<vector::InsertSlicesOp>(vectorDefOp);
     if (!insertSlicesOp)
-      return matchFailure();
+      return failure();
 
     // Get TupleOp operand of 'insertSlicesOp'.
     auto tupleOp = dyn_cast_or_null<vector::TupleOp>(
         insertSlicesOp.vectors().getDefiningOp());
     if (!tupleOp)
-      return matchFailure();
+      return failure();
 
     // Get 'sizes' and 'strides' parameters from InsertSlicesOp user.
     auto sourceTupleType = insertSlicesOp.getSourceTupleType();
     auto resultVectorType = insertSlicesOp.getResultVectorType();
     SmallVector<int64_t, 4> sizes;
     insertSlicesOp.getSizes(sizes);
     SmallVector<int64_t, 4> strides;
     insertSlicesOp.getStrides(strides);
auto createSlice = [&](unsigned index, ArrayRef<Value> sliceIndices) {
           xferWriteOp.permutation_map());
     };
     generateTransferOpSlices(memrefElementType, resultVectorType,
                              sourceTupleType, sizes, strides, indices, rewriter,
                              createSlice);
 
     // Erase old 'xferWriteOp'.
     rewriter.eraseOp(xferWriteOp);
-    return matchSuccess();
+    return success();
   }
 };
 
 /// Decomposes ShapeCastOp on tuple-of-vectors to multiple ShapeCastOps, each
 /// on vector types.
 struct ShapeCastOpDecomposer : public OpRewritePattern<vector::ShapeCastOp> {
   using OpRewritePattern<vector::ShapeCastOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(vector::ShapeCastOp shapeCastOp,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(vector::ShapeCastOp shapeCastOp,
     // Check if 'shapeCastOp' has tuple source/result type.
     auto sourceTupleType =
         shapeCastOp.source().getType().dyn_cast_or_null<TupleType>();
     auto resultTupleType =
         shapeCastOp.result().getType().dyn_cast_or_null<TupleType>();
     if (!sourceTupleType || !resultTupleType)
-      return matchFailure();
+      return failure();
     assert(sourceTupleType.size() == resultTupleType.size());
 
     // Create single-vector ShapeCastOp for each source tuple element.
     Location loc = shapeCastOp.getLoc();
     SmallVector<Value, 8> resultElements;
     resultElements.reserve(resultTupleType.size());
     for (unsigned i = 0, e = sourceTupleType.size(); i < e; ++i) {
       auto sourceElement = rewriter.create<vector::TupleGetOp>(
           loc, sourceTupleType.getType(i), shapeCastOp.source(),
           rewriter.getI64IntegerAttr(i));
       resultElements.push_back(rewriter.create<vector::ShapeCastOp>(
           loc, resultTupleType.getType(i), sourceElement));
     }
 
     // Replace 'shapeCastOp' with tuple of 'resultElements'.
     rewriter.replaceOpWithNewOp<vector::TupleOp>(shapeCastOp, resultTupleType,
                                                  resultElements);
-    return matchSuccess();
+    return success();
   }
 };
 
 /// ShapeCastOpFolder folds cancelling ShapeCastOps away.
 //
 // Example:
 //
 //  The following MLIR with cancelling ShapeCastOps:
 //
 //   %0 = source : vector<5x4x2xf32>
 //   %1 = shape_cast %0 : vector<5x4x2xf32> to vector<20x2xf32>
 //   %2 = shape_cast %1 : vector<20x2xf32> to vector<5x4x2xf32>
 //   %3 = user %2 : vector<5x4x2xf32>
 //
 //  Should canonicalize to the following:
 //
 //   %0 = source : vector<5x4x2xf32>
 //   %1 = user %0 : vector<5x4x2xf32>
 //
 struct ShapeCastOpFolder : public OpRewritePattern<vector::ShapeCastOp> {
   using OpRewritePattern<vector::ShapeCastOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(vector::ShapeCastOp shapeCastOp,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(vector::ShapeCastOp shapeCastOp,
     // Check if 'shapeCastOp' has vector source/result type.
     auto sourceVectorType =
         shapeCastOp.source().getType().dyn_cast_or_null<VectorType>();
     auto resultVectorType =
         shapeCastOp.result().getType().dyn_cast_or_null<VectorType>();
     if (!sourceVectorType || !resultVectorType)
-      return matchFailure();
+      return failure();
 
     // Check if shape cast op source operand is also a shape cast op.
     auto sourceShapeCastOp = dyn_cast_or_null<vector::ShapeCastOp>(
         shapeCastOp.source().getDefiningOp());
     if (!sourceShapeCastOp)
-      return matchFailure();
+      return failure();
     auto operandSourceVectorType =
         sourceShapeCastOp.source().getType().cast<VectorType>();
     auto operandResultVectorType =
         sourceShapeCastOp.result().getType().cast<VectorType>();
 
     // Check if shape cast operations invert each other.
     if (operandSourceVectorType != resultVectorType ||
         operandResultVectorType != sourceVectorType)
-      return matchFailure();
+      return failure();
 
     rewriter.replaceOp(shapeCastOp, sourceShapeCastOp.source());
-    return matchSuccess();
+    return success();
   }
 };
 
 // Patter rewrite which forward tuple elements to their users.
 // User(TupleGetOp(ExtractSlicesOp(InsertSlicesOp(TupleOp(Producer)))))
 //   -> User(Producer)
 struct TupleGetFolderOp : public OpRewritePattern<vector::TupleGetOp> {
   using OpRewritePattern<vector::TupleGetOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(vector::TupleGetOp tupleGetOp,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(vector::TupleGetOp tupleGetOp,
     // Return if 'tupleGetOp.vectors' arg was not defined by ExtractSlicesOp.
     auto extractSlicesOp = dyn_cast_or_null<vector::ExtractSlicesOp>(
         tupleGetOp.vectors().getDefiningOp());
     if (!extractSlicesOp)
-      return matchFailure();
+      return failure();
 
     // Return if 'extractSlicesOp.vector' arg was not defined by InsertSlicesOp.
     auto insertSlicesOp = dyn_cast_or_null<vector::InsertSlicesOp>(
         extractSlicesOp.vector().getDefiningOp());
     if (!insertSlicesOp)
-      return matchFailure();
+      return failure();
 
     // Return if 'insertSlicesOp.vectors' arg was not defined by TupleOp.
     auto tupleOp = dyn_cast_or_null<vector::TupleOp>(
         insertSlicesOp.vectors().getDefiningOp());
     if (!tupleOp)
-      return matchFailure();
+      return failure();
 
     // Forward Value from 'tupleOp' at 'tupleGetOp.index'.
     Value tupleValue = tupleOp.getOperand(tupleGetOp.getIndex());
     rewriter.replaceOp(tupleGetOp, tupleValue);
-    return matchSuccess();
+    return success();
   }
 };
 
 /// Progressive lowering of ExtractSlicesOp to tuple of StridedSliceOp.
 /// One:
 ///   %x = vector.extract_slices %0
 /// is replaced by:
 ///   %a = vector.strided_slice %0
 ///   %b = vector.strided_slice %0
 ///   ..
 ///   %x = vector.tuple %a, %b, ..
 class ExtractSlicesOpLowering
     : public OpRewritePattern<vector::ExtractSlicesOp> {
 public:
   using OpRewritePattern<vector::ExtractSlicesOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(vector::ExtractSlicesOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(vector::ExtractSlicesOp op,
     auto loc = op.getLoc();
 
     VectorType vectorType = op.getSourceVectorType();
     auto shape = vectorType.getShape();
 
     SmallVector<int64_t, 4> sizes;
     op.getSizes(sizes);
     SmallVector<int64_t, 4> strides;
for (int64_t i = 0; i < tupleSize; ++i) {
           computeElementOffsetsFromVectorSliceOffsets(sizes, vectorOffsets);
       auto sliceSizes = computeSliceSizes(shape, sizes, elementOffsets);
       // Insert in tuple.
       tupleValues[i] = rewriter.create<vector::StridedSliceOp>(
           loc, op.vector(), elementOffsets, sliceSizes, strides);
     }
 
     rewriter.replaceOpWithNewOp<vector::TupleOp>(op, tupleType, tupleValues);
-    return matchSuccess();
+    return success();
   }
 };
 
 /// Progressive lowering of InsertSlicesOp to series of InsertStridedSliceOp.
 /// One:
 ///   %x = vector.insert_slices %0
 /// is replaced by:
 ///   %r0 = vector.splat 0
 //    %t1 = vector.tuple_get %0, 0
 ///   %r1 = vector.insert_strided_slice %r0, %t1
 //    %t2 = vector.tuple_get %0, 1
 ///   %r2 = vector.insert_strided_slice %r1, %t2
 ///   ..
 ///   %x  = ..
 class InsertSlicesOpLowering : public OpRewritePattern<vector::InsertSlicesOp> {
 public:
   using OpRewritePattern<vector::InsertSlicesOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(vector::InsertSlicesOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(vector::InsertSlicesOp op,
     auto loc = op.getLoc();
 
     VectorType vectorType = op.getResultVectorType();
     auto shape = vectorType.getShape();
 
     SmallVector<int64_t, 4> sizes;
     op.getSizes(sizes);
     SmallVector<int64_t, 4> strides;
for (int64_t i = 0; i < tupleSize; ++i) {
       auto index = rewriter.getI64IntegerAttr(i);
       auto tupleGet = rewriter.create<vector::TupleGetOp>(
           loc, tupleType.getType(i), op.getOperand(), index);
       result = rewriter.create<vector::InsertStridedSliceOp>(
           loc, tupleGet, result, elementOffsets, strides);
     }
 
     rewriter.replaceOp(op, result);
-    return matchSuccess();
+    return success();
   }
 };
 
 /// Progressive lowering of OuterProductOp.
 /// One:
 ///   %x = vector.outerproduct %lhs, %rhs, %acc
 /// is replaced by:
 ///   %z = zero-result
 ///   %0 = vector.extract %lhs[0]
 ///   %1 = vector.broadcast %0
 ///   %2 = vector.extract %acc[0]
 ///   %3 = vector.fma %1, %arg1, %2
 ///   %4 = vector.insert %3, %z[0]
 ///   ..
 ///   %x = vector.insert %.., %..[N-1]
 ///
 class OuterProductOpLowering : public OpRewritePattern<vector::OuterProductOp> {
 public:
   using OpRewritePattern<vector::OuterProductOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(vector::OuterProductOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(vector::OuterProductOp op,
     auto loc = op.getLoc();
 
     VectorType rhsType = op.getOperandVectorTypeRHS();
     VectorType resType = op.getVectorType();
     Type eltType = resType.getElementType();
     Value acc = (op.acc().empty()) ? nullptr : op.acc()[0];
 
     Value zero = rewriter.create<ConstantOp>(loc, eltType,
                                              rewriter.getZeroAttr(eltType));
     Value result = rewriter.create<SplatOp>(loc, resType, zero);
     for (int64_t d = 0, e = resType.getDimSize(0); d < e; ++d) {
       auto pos = rewriter.getI64ArrayAttr(d);
       Value x = rewriter.create<vector::ExtractOp>(loc, eltType, op.lhs(), pos);
       Value b = rewriter.create<vector::BroadcastOp>(loc, rhsType, x);
       Value m;
       if (acc) {
         Value z = rewriter.create<vector::ExtractOp>(loc, rhsType, acc, pos);
         m = rewriter.create<vector::FMAOp>(loc, b, op.rhs(), z);
       } else {
         m = rewriter.create<MulFOp>(loc, b, op.rhs());
       }
       result = rewriter.create<vector::InsertOp>(loc, resType, m, result, pos);
     }
     rewriter.replaceOp(op, result);
-    return matchSuccess();
+    return success();
   }
 };
 
 /// Progressive lowering of ContractionOp.
 /// One:
 ///   %x = vector.contract with at least one free/batch dimension
 /// is replaced by:
 ///   %a = vector.contract with one less free/batch dimension
 ///   %b = vector.contract with one less free/batch dimension
 ///   ..
 ///   %x = combine %a %b ..
 /// until a pure contraction is reached (no free/batch dimensions),
 /// which is replaced by a fma/reduction op.
 ///
 /// TODO(ajcbik): break down into transpose/reshape/cast ops
 ///               when they become available to avoid code dup
 /// TODO(ajcbik): investigate lowering order impact on performance
 class ContractionOpLowering : public OpRewritePattern<vector::ContractionOp> {
 public:
   using OpRewritePattern<vector::ContractionOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(vector::ContractionOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(vector::ContractionOp op,
     // TODO(ajcbik): implement masks
     if (llvm::size(op.masks()) != 0)
-      return matchFailure();
+      return failure();
 
     // TODO(ntv, ajcbik): implement benefits, cost models, separate this out in
     // a new pattern.
     // TODO(ntv, fhahn): once row-major mode is available in LLVM's matrix
     // intrinsics, use that.
     if (lowerToLLVMMatrixIntrinsics &&
         isColumnMajorMatmul(op.indexing_maps())) {
       VectorType lhsType = op.getLhsType();
if (lowerToLLVMMatrixIntrinsics &&
       mul = rewriter.create<vector::ShapeCastOp>(op.getLoc(),
                                                  op.acc().getType(), mul);
       Type elementType = op.getLhsType().getElementType();
       assert(elementType.isIntOrFloat());
       if (elementType.isa<IntegerType>())
         rewriter.replaceOpWithNewOp<AddIOp>(op, op.acc(), mul);
       else
         rewriter.replaceOpWithNewOp<AddFOp>(op, op.acc(), mul);
-      return matchSuccess();
+      return success();
     }
 
     // Find first batch dimension in LHS/RHS, and lower when found.
     std::vector<std::pair<int64_t, int64_t>> batchDimMap = op.getBatchDimMap();
     if (!batchDimMap.empty()) {
       int64_t lhsIndex = batchDimMap[0].first;
       int64_t rhsIndex = batchDimMap[0].second;
       rewriter.replaceOp(op, lowerParallel(op, lhsIndex, rhsIndex, rewriter));
-      return matchSuccess();
+      return success();
     }
 
     // Collect contracting dimensions.
     std::vector<std::pair<int64_t, int64_t>> contractingDimMap =
         op.getContractingDimMap();
     DenseSet<int64_t> lhsContractingDimSet;
     DenseSet<int64_t> rhsContractingDimSet;
     for (auto &dimPair : contractingDimMap) {
       lhsContractingDimSet.insert(dimPair.first);
       rhsContractingDimSet.insert(dimPair.second);
     }
 
     // Find first free dimension in LHS, and lower when found.
     VectorType lhsType = op.getLhsType();
     for (int64_t lhsIndex = 0, e = lhsType.getRank(); lhsIndex < e;
          ++lhsIndex) {
       if (lhsContractingDimSet.count(lhsIndex) == 0) {
         rewriter.replaceOp(
             op, lowerParallel(op, lhsIndex, /*rhsIndex=*/-1, rewriter));
-        return matchSuccess();
+        return success();
       }
     }
 
     // Find first free dimension in RHS, and lower when found.
     VectorType rhsType = op.getRhsType();
     for (int64_t rhsIndex = 0, e = rhsType.getRank(); rhsIndex < e;
          ++rhsIndex) {
       if (rhsContractingDimSet.count(rhsIndex) == 0) {
         rewriter.replaceOp(
             op, lowerParallel(op, /*lhsIndex=*/-1, rhsIndex, rewriter));
-        return matchSuccess();
+        return success();
       }
     }
 
     // Lower the first remaining reduction dimension.
     if (!contractingDimMap.empty()) {
       rewriter.replaceOp(op, lowerReduction(op, rewriter));
-      return matchSuccess();
+      return success();
     }
 
-    return matchFailure();
+    return failure();
   }
 
 private:
   // Lower one parallel dimension.
   // TODO(ajcbik): consider reusing existing contract unrolling
   Value lowerParallel(vector::ContractionOp op, int64_t lhsIndex,
                       int64_t rhsIndex, PatternRewriter &rewriter) const {
     VectorType lhsType = op.getLhsType();

 /// This iterates over the most major dimension of the 2-D vector and performs
 /// rewrites into:
 ///   vector.extract from 2-D + vector.insert_strided_slice offset into 1-D
 class ShapeCastOp2DDownCastRewritePattern
     : public OpRewritePattern<vector::ShapeCastOp> {
 public:
   using OpRewritePattern<vector::ShapeCastOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(vector::ShapeCastOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(vector::ShapeCastOp op,
     auto sourceVectorType = op.getSourceVectorType();
     auto resultVectorType = op.getResultVectorType();
     if (sourceVectorType.getRank() != 2 || resultVectorType.getRank() != 1)
-      return matchFailure();
+      return failure();
 
     auto loc = op.getLoc();
     auto elemType = sourceVectorType.getElementType();
     Value zero = rewriter.create<ConstantOp>(loc, elemType,
                                              rewriter.getZeroAttr(elemType));
     Value desc = rewriter.create<SplatOp>(loc, resultVectorType, zero);
     unsigned mostMinorVectorSize = sourceVectorType.getShape()[1];
     for (int64_t i = 0, e = sourceVectorType.getShape().front(); i != e; ++i) {
       Value vec = rewriter.create<vector::ExtractOp>(loc, op.source(), i);
       desc = rewriter.create<vector::InsertStridedSliceOp>(
           loc, vec, desc,
           /*offsets=*/i * mostMinorVectorSize, /*strides=*/1);
     }
     rewriter.replaceOp(op, desc);
-    return matchSuccess();
+    return success();
   }
 };
 
 /// ShapeOp 1D -> 2D upcast serves the purpose of unflattening 2-D from 1-D
 /// vectors progressively on the way from targeting llvm.matrix intrinsics.
 /// This iterates over the most major dimension of the 2-D vector and performs
 /// rewrites into:
 ///   vector.strided_slice from 1-D + vector.insert into 2-D
 class ShapeCastOp2DUpCastRewritePattern
     : public OpRewritePattern<vector::ShapeCastOp> {
 public:
   using OpRewritePattern<vector::ShapeCastOp>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(vector::ShapeCastOp op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(vector::ShapeCastOp op,
     auto sourceVectorType = op.getSourceVectorType();
     auto resultVectorType = op.getResultVectorType();
     if (sourceVectorType.getRank() != 1 || resultVectorType.getRank() != 2)
-      return matchFailure();
+      return failure();
 
     auto loc = op.getLoc();
     auto elemType = sourceVectorType.getElementType();
     Value zero = rewriter.create<ConstantOp>(loc, elemType,
                                              rewriter.getZeroAttr(elemType));
     Value desc = rewriter.create<SplatOp>(loc, resultVectorType, zero);
     unsigned mostMinorVectorSize = resultVectorType.getShape()[1];
     for (int64_t i = 0, e = resultVectorType.getShape().front(); i != e; ++i) {
       Value vec = rewriter.create<vector::StridedSliceOp>(
           loc, op.source(), /*offsets=*/i * mostMinorVectorSize,
           /*sizes=*/mostMinorVectorSize,
           /*strides=*/1);
       desc = rewriter.create<vector::InsertOp>(loc, vec, desc, i);
     }
     rewriter.replaceOp(op, desc);
-    return matchSuccess();
+    return success();
   }
 };
 
 } // namespace
 
 // TODO(andydavis) Add pattern to rewrite ExtractSlices(ConstantMaskOp).
 // TODO(andydavis) Add this as DRR pattern.
 void mlir::vector::populateVectorToVectorTransformationPatterns(

mlir/lib/IR/PatternMatch.cpp

 // RewritePattern and PatternRewriter implementation
 //===----------------------------------------------------------------------===//
 
 void RewritePattern::rewrite(Operation *op, PatternRewriter &rewriter) const {
   llvm_unreachable("need to implement either matchAndRewrite or one of the "
                    "rewrite functions!");
 }
 
-PatternMatchResult RewritePattern::match(Operation *op) const {
+LogicalResult RewritePattern::match(Operation *op) const {
   llvm_unreachable("need to implement either match or matchAndRewrite!");
 }
 
 /// Patterns must specify the root operation name they match against, and can
 /// also specify the benefit of the pattern matching. They can also specify the
 /// names of operations that may be generated during a successful rewrite.
 RewritePattern::RewritePattern(StringRef rootName,
                                ArrayRef<StringRef> generatedNames,

mlir/lib/Quantizer/Transforms/RemoveInstrumentationPass.cpp

 };
 
 template <typename OpTy>
 class RemoveIdentityOpRewrite : public RewritePattern {
 public:
   RemoveIdentityOpRewrite(MLIRContext *context)
       : RewritePattern(OpTy::getOperationName(), 1, context) {}
 
-  PatternMatchResult matchAndRewrite(Operation *op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(Operation *op,
     assert(op->getNumOperands() == 1);
     assert(op->getNumResults() == 1);
 
     rewriter.replaceOp(op, op->getOperand(0));
-    return matchSuccess();
+    return success();
   }
 };
 
 } // end anonymous namespace
 
 void RemoveInstrumentationPass::runOnFunction() {
   OwningRewritePatternList patterns;
   auto func = getFunction();

mlir/lib/Transforms/DialectConversion.cpp

detail::ConversionPatternRewriterImpl &ConversionPatternRewriter::getImpl() {
   return *impl;
 }
 
 //===----------------------------------------------------------------------===//
 // Conversion Patterns
 //===----------------------------------------------------------------------===//
 
 /// Attempt to match and rewrite the IR root at the specified operation.
-PatternMatchResult
+LogicalResult
 ConversionPattern::matchAndRewrite(Operation *op,
                                    PatternRewriter &rewriter) const {
   SmallVector<Value, 4> operands;
   auto &dialectRewriter = static_cast<ConversionPatternRewriter &>(rewriter);
   dialectRewriter.getImpl().remapValues(op->getOperands(), operands);
   return matchAndRewrite(op, operands, dialectRewriter);
 }
 

 /// Create a default conversion pattern that rewrites the type signature of a
 /// FuncOp.
 namespace {
 struct FuncOpSignatureConversion : public OpConversionPattern<FuncOp> {
   FuncOpSignatureConversion(MLIRContext *ctx, TypeConverter &converter)
       : OpConversionPattern(ctx), converter(converter) {}
 
   /// Hook for derived classes to implement combined matching and rewriting.
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(FuncOp funcOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     FunctionType type = funcOp.getType();
 
     // Convert the original function arguments.
     TypeConverter::SignatureConversion result(type.getNumInputs());
     for (unsigned i = 0, e = type.getNumInputs(); i != e; ++i)
       if (failed(converter.convertSignatureArg(i, type.getInput(i), result)))
-        return matchFailure();
+        return failure();
 
     // Convert the original function results.
     SmallVector<Type, 1> convertedResults;
     if (failed(converter.convertTypes(type.getResults(), convertedResults)))
-      return matchFailure();
+      return failure();
 
     // Update the function signature in-place.
     rewriter.updateRootInPlace(funcOp, [&] {
       funcOp.setType(FunctionType::get(result.getConvertedTypes(),
                                        convertedResults, funcOp.getContext()));
       rewriter.applySignatureConversion(&funcOp.getBody(), result);
     });
-    return matchSuccess();
+    return success();
   }
 
   /// The type converter to use when rewriting the signature.
   TypeConverter &converter;
 };
 } // end anonymous namespace
 
 void mlir::populateFuncOpTypeConversionPattern(

mlir/test/lib/Dialect/SPIRV/TestAvailability.cpp

 namespace {
 /// A pass for testing SPIR-V op availability.
 struct ConvertToTargetEnv : public FunctionPass<ConvertToTargetEnv> {
   void runOnFunction() override;
 };
 
 struct ConvertToAtomCmpExchangeWeak : public RewritePattern {
   ConvertToAtomCmpExchangeWeak(MLIRContext *context);
-  PatternMatchResult matchAndRewrite(Operation *op,
-                                     PatternRewriter &rewriter) const override;
+  LogicalResult matchAndRewrite(Operation *op,
 };
 
 struct ConvertToBitReverse : public RewritePattern {
   ConvertToBitReverse(MLIRContext *context);
-  PatternMatchResult matchAndRewrite(Operation *op,
-                                     PatternRewriter &rewriter) const override;
+  LogicalResult matchAndRewrite(Operation *op,
 };
 
 struct ConvertToGroupNonUniformBallot : public RewritePattern {
   ConvertToGroupNonUniformBallot(MLIRContext *context);
-  PatternMatchResult matchAndRewrite(Operation *op,
-                                     PatternRewriter &rewriter) const override;
+  LogicalResult matchAndRewrite(Operation *op,
 };
 
 struct ConvertToModule : public RewritePattern {
   ConvertToModule(MLIRContext *context);
-  PatternMatchResult matchAndRewrite(Operation *op,
-                                     PatternRewriter &rewriter) const override;
+  LogicalResult matchAndRewrite(Operation *op,
 };
 
 struct ConvertToSubgroupBallot : public RewritePattern {
   ConvertToSubgroupBallot(MLIRContext *context);
-  PatternMatchResult matchAndRewrite(Operation *op,
-                                     PatternRewriter &rewriter) const override;
+  LogicalResult matchAndRewrite(Operation *op,
 };
 } // end anonymous namespace
 
 void ConvertToTargetEnv::runOnFunction() {
   MLIRContext *context = &getContext();
   FuncOp fn = getFunction();
 
   auto targetEnv = fn.getOperation()
void ConvertToTargetEnv::runOnFunction() {
   if (failed(applyPartialConversion(fn, *target, patterns)))
     return signalPassFailure();
 }
 
 ConvertToAtomCmpExchangeWeak::ConvertToAtomCmpExchangeWeak(MLIRContext *context)
     : RewritePattern("test.convert_to_atomic_compare_exchange_weak_op",
                      {"spv.AtomicCompareExchangeWeak"}, 1, context) {}
 
-PatternMatchResult
+LogicalResult
 ConvertToAtomCmpExchangeWeak::matchAndRewrite(Operation *op,
                                               PatternRewriter &rewriter) const {
   Value ptr = op->getOperand(0);
   Value value = op->getOperand(1);
   Value comparator = op->getOperand(2);
 
   // Create a spv.AtomicCompareExchangeWeak op with AtomicCounterMemory bits in
   // memory semantics to additionally require AtomicStorage capability.
   rewriter.replaceOpWithNewOp<spirv::AtomicCompareExchangeWeakOp>(
       op, value.getType(), ptr, spirv::Scope::Workgroup,
       spirv::MemorySemantics::AcquireRelease |
           spirv::MemorySemantics::AtomicCounterMemory,
       spirv::MemorySemantics::Acquire, value, comparator);
-  return matchSuccess();
+  return success();
 }
 
 ConvertToBitReverse::ConvertToBitReverse(MLIRContext *context)
     : RewritePattern("test.convert_to_bit_reverse_op", {"spv.BitReverse"}, 1,
                      context) {}
 
-PatternMatchResult
+LogicalResult
 ConvertToBitReverse::matchAndRewrite(Operation *op,
                                      PatternRewriter &rewriter) const {
   Value predicate = op->getOperand(0);
 
   rewriter.replaceOpWithNewOp<spirv::BitReverseOp>(
       op, op->getResult(0).getType(), predicate);
-  return matchSuccess();
+  return success();
 }
 
 ConvertToGroupNonUniformBallot::ConvertToGroupNonUniformBallot(
     MLIRContext *context)
     : RewritePattern("test.convert_to_group_non_uniform_ballot_op",
                      {"spv.GroupNonUniformBallot"}, 1, context) {}
 
-PatternMatchResult ConvertToGroupNonUniformBallot::matchAndRewrite(
+LogicalResult ConvertToGroupNonUniformBallot::matchAndRewrite(
     Operation *op, PatternRewriter &rewriter) const {
   Value predicate = op->getOperand(0);
 
   rewriter.replaceOpWithNewOp<spirv::GroupNonUniformBallotOp>(
       op, op->getResult(0).getType(), spirv::Scope::Workgroup, predicate);
-  return matchSuccess();
+  return success();
 }
 
 ConvertToModule::ConvertToModule(MLIRContext *context)
     : RewritePattern("test.convert_to_module_op", {"spv.module"}, 1, context) {}
 
-PatternMatchResult
+LogicalResult
 ConvertToModule::matchAndRewrite(Operation *op,
                                  PatternRewriter &rewriter) const {
   rewriter.replaceOpWithNewOp<spirv::ModuleOp>(
       op, spirv::AddressingModel::PhysicalStorageBuffer64,
       spirv::MemoryModel::Vulkan);
-  return matchSuccess();
+  return success();
 }
 
 ConvertToSubgroupBallot::ConvertToSubgroupBallot(MLIRContext *context)
     : RewritePattern("test.convert_to_subgroup_ballot_op",
                      {"spv.SubgroupBallotKHR"}, 1, context) {}
 
-PatternMatchResult
+LogicalResult
 ConvertToSubgroupBallot::matchAndRewrite(Operation *op,
                                          PatternRewriter &rewriter) const {
   Value predicate = op->getOperand(0);
 
   rewriter.replaceOpWithNewOp<spirv::SubgroupBallotKHROp>(
       op, op->getResult(0).getType(), predicate);
-  return matchSuccess();
+  return success();
 }
 
 namespace mlir {
 void registerConvertToTargetEnvPass() {
   PassRegistration<ConvertToTargetEnv> convertToTargetEnvPass(
       "test-spirv-target-env", "Test SPIR-V target environment");
 }
 } // namespace mlir

mlir/test/lib/TestDialect/TestDialect.cpp

 // Test removing op with inner ops.
 //===----------------------------------------------------------------------===//
 
 namespace {
 struct TestRemoveOpWithInnerOps
     : public OpRewritePattern<TestOpWithRegionPattern> {
   using OpRewritePattern<TestOpWithRegionPattern>::OpRewritePattern;
 
-  PatternMatchResult matchAndRewrite(TestOpWithRegionPattern op,
-                                     PatternRewriter &rewriter) const override {
+  LogicalResult matchAndRewrite(TestOpWithRegionPattern op,
     rewriter.eraseOp(op);
-    return matchSuccess();
+    return success();
   }
 };
 } // end anonymous namespace
 
 void TestOpWithRegionPattern::getCanonicalizationPatterns(
     OwningRewritePatternList &results, MLIRContext *context) {
   results.insert<TestRemoveOpWithInnerOps>(context);
 }

mlir/test/lib/TestDialect/TestPatterns.cpp

 // Region-Block Rewrite Testing
 
 /// This pattern is a simple pattern that inlines the first region of a given
 /// operation into the parent region.
 struct TestRegionRewriteBlockMovement : public ConversionPattern {
   TestRegionRewriteBlockMovement(MLIRContext *ctx)
       : ConversionPattern("test.region", 1, ctx) {}
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const final {
     // Inline this region into the parent region.
     auto &parentRegion = *op->getParentRegion();
     if (op->getAttr("legalizer.should_clone"))
       rewriter.cloneRegionBefore(op->getRegion(0), parentRegion,
                                  parentRegion.end());
     else
       rewriter.inlineRegionBefore(op->getRegion(0), parentRegion,
                                   parentRegion.end());
 
     // Drop this operation.
     rewriter.eraseOp(op);
-    return matchSuccess();
+    return success();
   }
 };
 /// This pattern is a simple pattern that generates a region containing an
 /// illegal operation.
 struct TestRegionRewriteUndo : public RewritePattern {
   TestRegionRewriteUndo(MLIRContext *ctx)
       : RewritePattern("test.region_builder", 1, ctx) {}
 
-  PatternMatchResult matchAndRewrite(Operation *op,
-                                     PatternRewriter &rewriter) const final {
+  LogicalResult matchAndRewrite(Operation *op,
     // Create the region operation with an entry block containing arguments.
     OperationState newRegion(op->getLoc(), "test.region");
     newRegion.addRegion();
     auto *regionOp = rewriter.createOperation(newRegion);
     auto *entryBlock = rewriter.createBlock(&regionOp->getRegion(0));
     entryBlock->addArgument(rewriter.getIntegerType(64));
 
     // Add an explicitly illegal operation to ensure the conversion fails.
     rewriter.create<ILLegalOpF>(op->getLoc(), rewriter.getIntegerType(32));
     rewriter.create<TestValidOp>(op->getLoc(), ArrayRef<Value>());
 
     // Drop this operation.
     rewriter.eraseOp(op);
-    return matchSuccess();
+    return success();
   }
 };
 
 //===----------------------------------------------------------------------===//
 // Type-Conversion Rewrite Testing
 
 /// This patterns erases a region operation that has had a type conversion.
 struct TestDropOpSignatureConversion : public ConversionPattern {
   TestDropOpSignatureConversion(MLIRContext *ctx, TypeConverter &converter)
       : ConversionPattern("test.drop_region_op", 1, ctx), converter(converter) {
   }
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     Region &region = op->getRegion(0);
     Block *entry = &region.front();
 
     // Convert the original entry arguments.
     TypeConverter::SignatureConversion result(entry->getNumArguments());
     for (unsigned i = 0, e = entry->getNumArguments(); i != e; ++i)
       if (failed(converter.convertSignatureArg(
               i, entry->getArgument(i).getType(), result)))
-        return matchFailure();
+        return failure();
 
     // Convert the region signature and just drop the operation.
     rewriter.applySignatureConversion(&region, result);
     rewriter.eraseOp(op);
-    return matchSuccess();
+    return success();
   }
 
   /// The type converter to use when rewriting the signature.
   TypeConverter &converter;
 };
 /// This pattern simply updates the operands of the given operation.
 struct TestPassthroughInvalidOp : public ConversionPattern {
   TestPassthroughInvalidOp(MLIRContext *ctx)
       : ConversionPattern("test.invalid", 1, ctx) {}
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const final {
     rewriter.replaceOpWithNewOp<TestValidOp>(op, llvm::None, operands,
                                              llvm::None);
-    return matchSuccess();
+    return success();
   }
 };
 /// This pattern handles the case of a split return value.
 struct TestSplitReturnType : public ConversionPattern {
   TestSplitReturnType(MLIRContext *ctx)
       : ConversionPattern("test.return", 1, ctx) {}
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const final {
     // Check for a return of F32.
     if (op->getNumOperands() != 1 || !op->getOperand(0).getType().isF32())
-      return matchFailure();
+      return failure();
 
     // Check if the first operation is a cast operation, if it is we use the
     // results directly.
     auto *defOp = operands[0].getDefiningOp();
     if (auto packerOp = llvm::dyn_cast_or_null<TestCastOp>(defOp)) {
       rewriter.replaceOpWithNewOp<TestReturnOp>(op, packerOp.getOperands());
-      return matchSuccess();
+      return success();
     }
 
     // Otherwise, fail to match.
-    return matchFailure();
+    return failure();
   }
 };
 
 //===----------------------------------------------------------------------===//
 // Multi-Level Type-Conversion Rewrite Testing
 struct TestChangeProducerTypeI32ToF32 : public ConversionPattern {
   TestChangeProducerTypeI32ToF32(MLIRContext *ctx)
       : ConversionPattern("test.type_producer", 1, ctx) {}
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const final {
     // If the type is I32, change the type to F32.
     if (!Type(*op->result_type_begin()).isSignlessInteger(32))
-      return matchFailure();
+      return failure();
     rewriter.replaceOpWithNewOp<TestTypeProducerOp>(op, rewriter.getF32Type());
-    return matchSuccess();
+    return success();
   }
 };
 struct TestChangeProducerTypeF32ToF64 : public ConversionPattern {
   TestChangeProducerTypeF32ToF64(MLIRContext *ctx)
       : ConversionPattern("test.type_producer", 1, ctx) {}
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const final {
     // If the type is F32, change the type to F64.
     if (!Type(*op->result_type_begin()).isF32())
       return rewriter.notifyMatchFailure(op, "expected single f32 operand");
     rewriter.replaceOpWithNewOp<TestTypeProducerOp>(op, rewriter.getF64Type());
-    return matchSuccess();
+    return success();
   }
 };
 struct TestChangeProducerTypeF32ToInvalid : public ConversionPattern {
   TestChangeProducerTypeF32ToInvalid(MLIRContext *ctx)
       : ConversionPattern("test.type_producer", 10, ctx) {}
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const final {
     // Always convert to B16, even though it is not a legal type. This tests
     // that values are unmapped correctly.
     rewriter.replaceOpWithNewOp<TestTypeProducerOp>(op, rewriter.getBF16Type());
-    return matchSuccess();
+    return success();
   }
 };
 struct TestUpdateConsumerType : public ConversionPattern {
   TestUpdateConsumerType(MLIRContext *ctx)
       : ConversionPattern("test.type_consumer", 1, ctx) {}
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const final {
     // Verify that the incoming operand has been successfully remapped to F64.
     if (!operands[0].getType().isF64())
-      return matchFailure();
+      return failure();
     rewriter.replaceOpWithNewOp<TestTypeConsumerOp>(op, operands[0]);
-    return matchSuccess();
+    return success();
   }
 };
 
 //===----------------------------------------------------------------------===//
 // Non-Root Replacement Rewrite Testing
 /// This pattern generates an invalid operation, but replaces it before the
 /// pattern is finished. This checks that we don't need to legalize the
 /// temporary op.
 struct TestNonRootReplacement : public RewritePattern {
   TestNonRootReplacement(MLIRContext *ctx)
       : RewritePattern("test.replace_non_root", 1, ctx) {}
 
-  PatternMatchResult matchAndRewrite(Operation *op,
-                                     PatternRewriter &rewriter) const final {
+  LogicalResult matchAndRewrite(Operation *op,
     auto resultType = *op->result_type_begin();
     auto illegalOp = rewriter.create<ILLegalOpF>(op->getLoc(), resultType);
     auto legalOp = rewriter.create<LegalOpB>(op->getLoc(), resultType);
 
     rewriter.replaceOp(illegalOp, {legalOp});
     rewriter.replaceOp(op, {illegalOp});
-    return matchSuccess();
+    return success();
   }
 };
 } // namespace
 
 namespace {
 struct TestTypeConverter : public TypeConverter {
   using TypeConverter::TypeConverter;
   TestTypeConverter() { addConversion(convertType); }

 /// Example:
 ///   %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;
 
-  PatternMatchResult
+  LogicalResult
   matchAndRewrite(OneVResOneVOperandOp1 op, ArrayRef<Value> operands,
                   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'.
     remappedOperands.push_back(rewriter.getRemappedValue(origOp));
     remappedOperands.push_back(rewriter.getRemappedValue(origOp));
 
     rewriter.replaceOpWithNewOp<OneVResOneVOperandOp1>(op, op.getResultTypes(),
                                                        remappedOperands);
-    return matchSuccess();
+    return success();
   }
 };
 
 struct TestRemappedValue : public mlir::FunctionPass<TestRemappedValue> {
   void runOnFunction() override {
     mlir::OwningRewritePatternList patterns;
     patterns.insert<OneVResOneVOperandOp1Converter>(&getContext());
 

mlir/tools/mlir-tblgen/RewriterGen.cpp

void PatternEmitter::emitOpMatch(DagNode tree, int depth) {
 
   int indent = 4 + 2 * depth;
   os.indent(indent) << formatv(
       "auto castedOp{0} = dyn_cast_or_null<{1}>(op{0}); (void)castedOp{0};\n",
       depth, op.getQualCppClassName());
   // Skip the operand matching at depth 0 as the pattern rewriter already does.
   if (depth != 0) {
     // Skip if there is no defining operation (e.g., arguments to function).
-    os.indent(indent) << formatv("if (!castedOp{0}) return matchFailure();\n",
+    os.indent(indent) << formatv("if (!castedOp{0}) return failure();\n",
                                  depth);
   }
   if (tree.getNumArgs() != op.getNumArgs()) {
     PrintFatalError(loc, formatv("op '{0}' argument number mismatch: {1} in "
                                  "pattern vs. {2} in definition",
                                  op.getOperationName(), tree.getNumArgs(),
                                  op.getNumArgs()));
   }
if (operand->constraint != matcher.getAsConstraint()) {
         PrintFatalError(loc, error);
       }
       auto self =
           formatv("(*castedOp{0}.getODSOperands({1}).begin()).getType()", depth,
                   argIndex);
       os.indent(indent) << "if (!("
                         << std::string(tgfmt(matcher.getConditionTemplate(),
                                              &fmtCtx.withSelf(self)))
-                        << ")) return matchFailure();\n";
+                        << ")) return failure();\n";
     }
   }
 
   // Capture the value
   auto name = tree.getArgName(argIndex);
   // `$_` is a special symbol to ignore op argument matching.
   if (!name.empty() && name != "_") {
     // We need to subtract the number of attributes before this operand to get
os.indent(indent) << "if (!tblgen_attr) tblgen_attr = "
                       << std::string(tgfmt(attr.getConstBuilderTemplate(),
                                            &fmtCtx, attr.getDefaultValue()))
                       << ";\n";
   } else if (attr.isOptional()) {
     // For a missing attribute that is optional according to definition, we
     // should just capture a mlir::Attribute() to signal the missing state.
     // That is precisely what getAttr() returns on missing attributes.
   } else {
-    os.indent(indent) << "if (!tblgen_attr) return matchFailure();\n";
+    os.indent(indent) << "if (!tblgen_attr) return failure();\n";
   }
 
   auto matcher = tree.getArgAsLeaf(argIndex);
   if (!matcher.isUnspecified()) {
     if (!matcher.isAttrMatcher()) {
       PrintFatalError(
           loc, formatv("the {1}-th argument of op '{0}' should be an attribute",
                        op.getOperationName(), argIndex + 1));
     }
 
     // If a constraint is specified, we need to generate C++ statements to
     // check the constraint.
     os.indent(indent) << "if (!("
                       << std::string(tgfmt(matcher.getConditionTemplate(),
                                            &fmtCtx.withSelf("tblgen_attr")))
-                      << ")) return matchFailure();\n";
+                      << ")) return failure();\n";
   }
 
   // Capture the value
   auto name = tree.getArgName(argIndex);
   // `$_` is a special symbol to ignore op argument matching.
   if (!name.empty() && name != "_") {
     os.indent(indent) << formatv("{0} = tblgen_attr;\n", name);
   }
 
   indent -= 2;
   os.indent(indent) << "}\n";
 }
 
 void PatternEmitter::emitMatchLogic(DagNode tree) {
   LLVM_DEBUG(llvm::dbgs() << "--- start emitting match logic ---\n");
   emitOpMatch(tree, 0);
 
   for (auto &appliedConstraint : pattern.getConstraints()) {
     auto &constraint = appliedConstraint.constraint;
     auto &entities = appliedConstraint.entities;
 
     auto condition = constraint.getConditionTemplate();
-    auto cmd = "if (!({0})) return matchFailure();\n";
+    auto cmd = "if (!({0})) return failure();\n";
 
     if (isa<TypeConstraint>(constraint)) {
       auto self = formatv("({0}.getType())",
                           symbolInfoMap.getValueAndRangeUse(entities.front()));
       os.indent(4) << formatv(cmd,
                               tgfmt(condition, &fmtCtx.withSelf(self.str())));
     } else if (isa<AttrConstraint>(constraint)) {
       PrintFatalError(
void PatternEmitter::emit(StringRef rewriteName) {
   });
   interleaveComma(sortedResultOps, os, [&](const Operator *op) {
     os << '"' << op->getOperationName() << '"';
   });
   os << formatv(R"(}, {0}, context) {{})", pattern.getBenefit()) << "\n";
 
   // Emit matchAndRewrite() function.
   os << R"(
-  PatternMatchResult matchAndRewrite(Operation *op0,
+  LogicalResult matchAndRewrite(Operation *op0,
                                      PatternRewriter &rewriter) const override {
 )";
 
   // Register all symbols bound in the source pattern.
   pattern.collectSourcePatternBoundSymbols(symbolInfoMap);
 
   LLVM_DEBUG(
       llvm::dbgs() << "start creating local variables for capturing matches\n");
)";
   os.indent(4) << "// Match\n";
   os.indent(4) << "tblgen_ops[0] = op0;\n";
   emitMatchLogic(sourceTree);
   os << "\n";
 
   os.indent(4) << "// Rewrite\n";
   emitRewriteLogic();
 
-  os.indent(4) << "return matchSuccess();\n";
+  os.indent(4) << "return success();\n";
   os << "  };\n";
   os << "};\n";
 }
 
 void PatternEmitter::emitRewriteLogic() {
   LLVM_DEBUG(llvm::dbgs() << "--- start emitting rewrite logic ---\n");
   const Operator &rootOp = pattern.getSourceRootOp();
   int numExpectedResults = rootOp.getNumResults();