[mlir][transform] Use separate ops instead of PatternRegistry

[springerm]

Summary

• Remove transform::PatternRegistry.
• Add a new op for each currently registered pattern set.
• Change names of vector dialect pattern selector ops, so that they are consistent with the remaining code base.
• Remove redundant transform.vector.extract_address_computations op.

[ftynse]

mlir/test/Dialect/MemRef/extract-address-computations.mlir

57

Nit: something went wrong with formatting, the trailing } belongs to the next line.

mlir/include/mlir/Dialect/Linalg/TransformOps/LinalgTransformOps.td

 // This is roughly similar to OpFoldResult assuming the handle produces a single
 // value in the payload IR.
 def TransformParamTypeOrAnyHandle : Type<
     Or<[TransformHandleTypeInterface.predicate,
         Transform_ParamType.predicate]>,
     "transform 'param' type or any handle type">;
 
 //===----------------------------------------------------------------------===//
+// Apply...PatternsOp
+//===----------------------------------------------------------------------===//
+
+def ApplyEraseUnnecessaryInputsPatternsOp : Op<Transform_Dialect,
+    "apply_patterns.linalg.erase_unnecessary_inputs",
+    [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
+  let description = [{
+    Collects patterns that promote inputs to outputs and remove unused inputs of
+    `linalg.generic` ops.
+  }];
+
+  let assemblyFormat = "attr-dict";
+}
+
+def ApplyFoldUnitExtentDimsViaReshapesPatternsOp : Op<Transform_Dialect,
+    "apply_patterns.linalg.fold_unit_extent_dims_via_reshapes",
+    [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
+  let description = [{
+    Collects patterns to fold unit-extent dimensions in operands/results of
+    linalg ops on tensors via reassociative reshape ops.
+  }];
+
+  let assemblyFormat = "attr-dict";
+}
+
+def ApplyFoldUnitExtentDimsViaSlicesPatternsOp : Op<Transform_Dialect,
+    "apply_patterns.linalg.fold_unit_extent_dims_via_slices",
+    [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
+  let description = [{
+    Collects patterns to fold unit-extent dimensions in operands/results of
+    linalg ops on tensors via rank-reducing slices.
+  }];
+
+  let assemblyFormat = "attr-dict";
+}
+
+def ApplyTilingCanonicalizationPatternsOp : Op<Transform_Dialect,
+    "apply_patterns.linalg.tiling_canonicalization",
+    [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
+  let description = [{
+    Collects canonicalization patterns relevant to apply after tiling patterns.
+  }];
+
+  let assemblyFormat = "attr-dict";
+}
+
+//===----------------------------------------------------------------------===//
 // BufferizeToAllocationOp
 //===----------------------------------------------------------------------===//
 
 def BufferizeToAllocationOp : Op<Transform_Dialect,
     "structured.bufferize_to_allocation",
     [DeclareOpInterfaceMethods<TransformOpInterface>,
      DeclareOpInterfaceMethods<MemoryEffectsOpInterface>]> {
   let description = [{

mlir/include/mlir/Dialect/MemRef/TransformOps/MemRefTransformOps.td

 #define MEMREF_TRANSFORM_OPS
 
 include "mlir/Dialect/Transform/IR/TransformDialect.td"
 include "mlir/Dialect/Transform/IR/TransformInterfaces.td"
 include "mlir/Dialect/Transform/IR/TransformTypes.td"
 include "mlir/Interfaces/SideEffectInterfaces.td"
 include "mlir/IR/OpBase.td"
 
+def ApplyExpandOpsPatternsOp : Op<Transform_Dialect,
+    "apply_patterns.memref.expand_ops",
+    [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
+  let description = [{
+    Collects patterns to rewrite ops within the memref dialect.
+
+    - Converts `atomic_rmw` that cannot be lowered to a simple atomic op with
+      AtomicRMWOpLowering pattern, e.g. with "minf" or "maxf" attributes, to
+      `memref.generic_atomic_rmw` with the expanded code.
+    - Converts `memref.reshape` that has a target shape of a statically-known
+      size to `memref.reinterpret_cast`.
+  }];
+
+  let assemblyFormat = "attr-dict";
+}
+
+def ApplyExpandStridedMetadataPatternsOp : Op<Transform_Dialect,
+    "apply_patterns.memref.expand_strided_metadata",
+    [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
+  let description = [{
+    Collects patterns for expanding memref operations that modify the metadata
+    (sizes, offset, strides) of a memref into easier to analyze constructs.
+  }];
+
+  let assemblyFormat = "attr-dict";
+}
+
+def ApplyExtractAddressComputationsPatternsOp : Op<Transform_Dialect,
+    "apply_patterns.memref.extract_address_computations",
+    [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
+  let description = [{
+    Collects patterns for extracting address computations from operations
+    with memory accesses such that these memory accesses use only a base
+    pointer.
+
+    For instance,
+    ```mlir
+    memref.load %base[%off0, ...]
+    ```
+
+    Will be rewritten in:
+    ```mlir
+    %new_base = memref.subview %base[%off0,...][1,...][1,...]
+    memref.load %new_base[%c0,...]
+    ```
+  }];
+
+  let assemblyFormat = "attr-dict";
+}
+
+def ApplyFoldMemrefAliasOpsPatternsOp : Op<Transform_Dialect,
+    "apply_patterns.memref.fold_memref_alias_ops",
+    [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
+  let description = [{
+    Collects patterns for folding memref aliasing ops (memref.subview) into
+    consumer load/store ops (affine.load, memref.load, nvgpu.ldmatrix,
+    vector.load, vector.transfer_read, affine.store, memref.store, etc.) and
+    other ops (e.g., memref.subview).
+  }];
+
+  let assemblyFormat = "attr-dict";
+}
+
+def ApplyResolveRankedShapedTypeResultDimsPatternsOp : Op<Transform_Dialect,
+    "apply_patterns.memref.resolve_ranked_shaped_type_result_dims",
+    [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
+  let description = [{
+    Collects patterns that resolve `memref.dim` operations with values that are
+    defined by operations that implement the `ReifyRankedShapedTypeOpInterface`,
+    in terms of shapes of its input operands.
+  }];
+
+  let assemblyFormat = "attr-dict";
+}
+
 def Transform_MemRefAllocOp : Transform_ConcreteOpType<"memref.alloc">;
 
 def MemRefMultiBufferOp : Op<Transform_Dialect, "memref.multibuffer",
     [FunctionalStyleTransformOpTrait, MemoryEffectsOpInterface,
      DeclareOpInterfaceMethods<TransformOpInterface>]> {
   let summary = "Multibuffers an allocation";
   let description = [{
      Transformation to do multi-buffering/array expansion to remove
let arguments =
            UnitAttr:$skip_analysis);
 
   let results = (outs TransformHandleTypeInterface:$transformed);
 
   let assemblyFormat =
     "$target attr-dict `:` functional-type(operands, results)";
 }
 
-def MemRefExtractAddressComputationsOp :
-    Op<Transform_Dialect, "memref.extract_address_computations",
-       [FunctionalStyleTransformOpTrait, MemoryEffectsOpInterface,
-        TransformEachOpTrait, TransformOpInterface]> {
-  let summary = "Extract address computations from memory accesses";
-  let description = [{
-    Transformation that extracts address computations from instructions
-    with memory accesses such that these memory accesses use only a base
-    pointer.
-
-    For instance,
-    ```mlir
-    memref.load %base[%off0, ...]
-    ```
-
-    Will be rewritten in:
-    ```mlir
-    %new_base = memref.subview %base[%off0,...][1,...][1,...]
-    memref.load %new_base[%c0,...]
-    ```
-
-    Note: The current implementation requires that the input operation
-    is "isolated from above".
-
-    #### Return modes
-
-    This operation produces `definiteFailure` if the extraction fails for any
-    reason.
-    The operation always returns the handle to the target op that is expected
-    to be isolated from above.
-  }];
-
-  let arguments = (ins TransformHandleTypeInterface:$target);
-  let results = (outs TransformHandleTypeInterface:$transformed);
-
-  let assemblyFormat = "$target attr-dict `:` functional-type(operands, results)";
-
-  let extraClassDeclaration = [{
-    ::mlir::DiagnosedSilenceableFailure applyToOne(
-        ::mlir::Operation *target,
-        ::mlir::transform::ApplyToEachResultList &transformResults,
-        ::mlir::transform::TransformState &state);
-  }];
-}
-
 def MemRefMakeLoopIndependentOp
     : Op<Transform_Dialect, "memref.make_loop_independent",
          [FunctionalStyleTransformOpTrait, MemoryEffectsOpInterface,
           TransformOpInterface, TransformEachOpTrait]> {
   let description = [{
     Rewrite the targeted ops such that their index-typed operands no longer
     depend on any loop induction variable of the `num_loop` enclosing `scf.for`
     loops. I.e., compute an upper bound that is independent of any such loop IV

mlir/include/mlir/Dialect/SCF/TransformOps/SCFTransformOps.td

 #define SCF_TRANSFORM_OPS
 
 include "mlir/Dialect/Transform/IR/TransformDialect.td"
 include "mlir/Dialect/Transform/IR/TransformInterfaces.td"
 include "mlir/Dialect/Transform/IR/TransformTypes.td"
 include "mlir/Interfaces/SideEffectInterfaces.td"
 include "mlir/IR/OpBase.td"
 
+def ApplyForLoopCanonicalizationPatternsOp : Op<Transform_Dialect,
+    "apply_patterns.scf.for_loop_canonicalization",
+    [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
+  let description = [{
+    Collects patterns for canonicalizing operations inside SCF loop bodies.
+    At the moment, only affine.min/max computations with iteration variables,
+    loop bounds and loop steps are canonicalized.
+  }];
+
+  let assemblyFormat = "attr-dict";
+}
+
 def Transform_ScfForOp : Transform_ConcreteOpType<"scf.for">;
 
 def GetParentForOp : Op<Transform_Dialect, "loop.get_parent_for",
     [NavigationTransformOpTrait, MemoryEffectsOpInterface,
      DeclareOpInterfaceMethods<TransformOpInterface>]> {
   let summary = "Gets a handle to the parent 'for' loop of the given operation";
   let description = [{
     Produces a handle to the n-th (default 1) parent `scf.for` or `affine.for`

mlir/include/mlir/Dialect/Tensor/TransformOps/TensorTransformOps.h

 //===- TensorTransformOps.h - Tensor transformation ops ---------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef MLIR_DIALECT_TENSOR_TRANSFORMOPS_TENSORTRANSFORMOPS_H
 #define MLIR_DIALECT_TENSOR_TRANSFORMOPS_TENSORTRANSFORMOPS_H
 
-#include "mlir/Dialect/Transform/IR/TransformOps.h"
+#include "mlir/Dialect/Transform/IR/TransformInterfaces.h"
 #include "mlir/Dialect/Transform/IR/TransformTypes.h"
 #include "mlir/IR/OpImplementation.h"
 #include "mlir/IR/PatternMatch.h"
 
 namespace mlir {
 class DialectRegistry;
 
 namespace tensor {

mlir/include/mlir/Dialect/Tensor/TransformOps/TensorTransformOps.td

 #define TENSOR_TRANSFORM_OPS
 
 include "mlir/Dialect/Transform/IR/TransformDialect.td"
 include "mlir/Dialect/Transform/IR/TransformInterfaces.td"
 include "mlir/Dialect/Transform/IR/TransformTypes.td"
 include "mlir/Interfaces/SideEffectInterfaces.td"
 include "mlir/IR/OpBase.td"
 
+def ApplyDropRedundantInsertSliceRankExpansionPatternsOp : Op<Transform_Dialect,
+    "apply_patterns.tensor.drop_redundant_insert_slice_rank_expansion",
+    [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
+  let description = [{
+    Indicates that redundant tensor.insert_slice rank reductions should be
+    dropped. E.g., cases where a tensor.extract_slice rank reduction immediately
+    follows an inverse tensor.insert_slice rank expansion.
+  }];
+
+  let assemblyFormat = "attr-dict";
+}
+
+def ApplyFoldTensorEmptyPatternsOp : Op<Transform_Dialect,
+    "apply_patterns.tensor.fold_tensor_empty",
+    [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
+  let description = [{
+    Indicates that reassociative reshapes (tensor.collapse_shape /
+    tensor.expand_shape) should be folded with inverse rank expansions / rank
+    reductions (via tensor.insert_slice / tensor.extract_slice).
+  }];
+
+  let assemblyFormat = "attr-dict";
+}
+def ApplyFoldIntoPackAndUnpackPatternsOp : Op<Transform_Dialect,
+    "apply_patterns.tensor.fold_into_pack_and_unpack",
+    [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
+  let description = [{
+    Indicates that operations like tensor.pad and tensor.extract_slice should
+    be folded into tensor.pack and tensor.unpack operations, respectively.
+  }];
+
+  let assemblyFormat = "attr-dict";
+}
+
+def ApplyFoldTensorSubsetOpsPatternsOp : Op<Transform_Dialect,
+    "apply_patterns.tensor.fold_tensor_subset_ops",
+    [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
+  let description = [{
+    Indicates that tensor.empty should be folded with tensor.extract_slice,
+    tensor.expand_shape and tensor.collapse_shape.
+  }];
+
+  let assemblyFormat = "attr-dict";
+}
+
+def ApplyMergeConsecutiveInsertExtractSlicePatternsOp : Op<Transform_Dialect,
+    "apply_patterns.tensor.merge_consecutive_insert_extract_slice",
+    [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
+  let description = [{
+    Indicates that consecutive tensor.extract_slice/tensor.insert_slice ops
+    should be merged into a single op. These patterns are not canonicalizations
+    because the bufferization is sensitive to IR structure.
+  }];
+
+  let assemblyFormat = "attr-dict";
+}
+
+def ApplyReassociativeReshapeFoldingPatternsOp : Op<Transform_Dialect,
+    "apply_patterns.tensor.reassociative_reshape_folding",
+    [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
+  let description = [{
+    Indicates that reassociative reshapes (tensor.collapse_shape /
+    tensor.expand_shape) should be folded with inverse rank expansions / rank
+    reductions (via tensor.insert_slice / tensor.extract_slice).
+  }];
+
+  let assemblyFormat = "attr-dict";
+}
+
 def Transform_TensorPadOp : Transform_ConcreteOpType<"tensor.pad">;
 
 def MakeLoopIndependentOp
     : Op<Transform_Dialect, "tensor.make_loop_independent",
          [FunctionalStyleTransformOpTrait, MemoryEffectsOpInterface,
           TransformOpInterface, TransformEachOpTrait]> {
   let description = [{
     Rewrite the targeted ops such that their index-typed operands no longer

mlir/include/mlir/Dialect/Tensor/Transforms/Transforms.h

 // Populate functions.
 //===----------------------------------------------------------------------===//
 
 /// Appends patterns for folding tensor aliasing ops into consumer load/store
 /// ops into `patterns`.
 void populateFoldTensorSubsetOpPatterns(RewritePatternSet &patterns);
 
 /// Collects patterns to merge consecutive tensor.insert_slice/extract_slice
-/// into one. These patterns are in in this separate entry point because the
-/// bufferization is sensitive over IR structure, particularly those
+/// into one. These patterns are in this separate entry point because the
+/// bufferization is sensitive to IR structure, particularly those
 /// tensor.extract_slice and tensor.insert_slice ops for creating the slices.
 void populateMergeConsecutiveInsertExtractSlicePatterns(
     RewritePatternSet &patterns);
 
 /// Populates `patterns` with patterns that drop redundant tensor.insert_slice
 /// rank expansions.
 void populateDropRedundantInsertSliceRankExpansionPatterns(
     RewritePatternSet &patterns);
 
 /// Populates `patterns` with patterns that fold `tensor.expand_shape` and
 /// `tensor.collapse_shape` into other ops.
 void populateReassociativeReshapeFoldingPatterns(RewritePatternSet &patterns);
 
 /// Populates `patterns` with patterns that fold tensor.empty with
-/// tensor.[extract_slice|cast|expand_shape|collapse_shape].
+/// tensor.[extract_slice|expand_shape|collapse_shape].
 void populateFoldTensorEmptyPatterns(RewritePatternSet &patterns);
 
 /// Populates `patterns` with patterns that fold operations like `tensor.pad`
 /// and `tensor.extract_slice` into `tensor.pack` and `tensor.unpack` operations
 /// respectively.
 void populateFoldIntoPackAndUnpackPatterns(RewritePatternSet &patterns);
 
 //===----------------------------------------------------------------------===//

mlir/include/mlir/Dialect/Transform/IR/TransformOps.h

 #include "mlir/IR/PatternMatch.h"
 #include "mlir/IR/SymbolTable.h"
 #include "mlir/Interfaces/CallInterfaces.h"
 #include "mlir/Interfaces/CastInterfaces.h"
 #include "mlir/Interfaces/ControlFlowInterfaces.h"
 
 namespace mlir {
 namespace transform {
-class ApplyPatternsOp;
 
 enum class FailurePropagationMode : uint32_t;
 class FailurePropagationModeAttr;
 
 /// A builder function that populates the body of a SequenceOp.
 using SequenceBodyBuilderFn = ::llvm::function_ref<void(
     ::mlir::OpBuilder &, ::mlir::Location, ::mlir::BlockArgument)>;
 using SequenceBodyBuilderArgsFn =
private:
   /// The error state of this listener. "Success" indicates that no error
   /// happened so far.
   DiagnosedSilenceableFailure status = DiagnosedSilenceableFailure::success();
 
   /// The number of errors that have been encountered.
   int64_t errorCounter = 0;
 };
 
-/// The PatternRegistry stores callbacks to functions that populate a
-/// `RewritePatternSet`. Registered patterns can be applied with the
-/// "transform.apply_patterns" op.
-class PatternRegistry : public TransformDialectData<PatternRegistry> {
-public:
-  PatternRegistry(MLIRContext *ctx) : TransformDialectData(ctx), builder(ctx) {}
-
-  /// A function that populates a `RewritePatternSet`.
-  using PopulatePatternsFn = std::function<void(RewritePatternSet &)>;
-  /// A function that populates a `RewritePatternSet` with a specified benefit.
-  using PopulatePatternsWithBenefitFn =
-      std::function<void(RewritePatternSet &, PatternBenefit)>;
-
-  /// Registers patterns with the specified identifier. The identifier should
-  /// be prefixed with the dialect to which the patterns belong.
-  void registerPatterns(StringRef identifier, PopulatePatternsFn &&fn);
-
-  /// Registers patterns with the specified identifier. The identifier should
-  /// be prefixed with the dialect to which the patterns belong. The pattern
-  /// benefit is currently ignored.
-  void registerPatterns(StringRef identifier,
-                        PopulatePatternsWithBenefitFn &&fn);
-
-protected:
-  friend class ApplyPatternsOp;
-
-  /// Returns "true" if patterns are registered with the specified identifier.
-  bool hasPatterns(StringAttr identifier) const;
-
-  /// Populates the given pattern set with the specified patterns.
-  void populatePatterns(StringAttr identifier,
-                        RewritePatternSet &patternSet) const;
-
-private:
-  /// A builder for creating StringAttrs.
-  Builder builder;
-
-  DenseMap<StringAttr, PopulatePatternsFn> patterns;
-};
-
 } // namespace transform
 } // namespace mlir
 
-MLIR_DECLARE_EXPLICIT_TYPE_ID(mlir::transform::PatternRegistry)
-
 #define GET_OP_CLASSES
 #include "mlir/Dialect/Transform/IR/TransformOps.h.inc"
 
 #endif // MLIR_DIALECT_TRANSFORM_IR_TRANSFORMOPS_H

mlir/include/mlir/Dialect/Transform/IR/TransformOps.td

DeclareOpInterfaceMethods<MemoryEffectsOpInterface>]
         # GraphRegionNoTerminator.traits> {
   let summary = "Greedily applies patterns to the body of the targeted op";
   let description = [{
     This transform greedily applies the specified patterns to the body of the
     targeted op until a fixpoint was reached. Patterns are not applied to the
     targeted op itself.
 
     The patterns that should be applied are specified in the graph region of
-    this op. They must implement the `PatternDescriptorOpInterface`.
-
-    (Deprecated) In addition, patterns that were registered in the transform
-    dialect's `PatternRegistry` are available. "canonicalization" is a special
-    set of patterns that refers to all canonicalization patterns of all loaded
-    dialects.
+    this op. They must implement the `PatternDescriptorOpInterface`. The order
+    in which patterns are applied is unspecified; i.e., the ordering of ops in
+    the region of this op is irrelevant.
 
     This transform only reads the target handle and modifies the payload. If a
     pattern erases or replaces a tracked op, the mapping is updated accordingly.
 
     Only replacements via `RewriterBase::replaceOp` or `replaceOpWithNewOp` are
     considered "payload op replacements". Furthermore, only if the replacement
     values are defined by the same op and that op has the same type as the
     original op, the mapping is updated. Otherwise, this transform fails
     silently unless `fail_on_payload_replacement_not_found` is set to "false".
     More details can be found at the documentation site of `TrackingListener`.
 
     This transform also fails silently if the pattern application did not
     converge within the default number of iterations/rewrites of the greedy
     pattern rewrite driver.
   }];
 
   let arguments = (ins
-    TransformHandleTypeInterface:$target, ArrayAttr:$patterns,
+    TransformHandleTypeInterface:$target,
     DefaultValuedAttr<BoolAttr, "true">:$fail_on_payload_replacement_not_found);
   let results = (outs);
   let regions = (region MaxSizedRegion<1>:$region);
 
-  let assemblyFormat = "$patterns `to` $target $region attr-dict `:` type($target)";
+  let assemblyFormat = "`to` $target $region attr-dict `:` type($target)";
   let hasVerifier = 1;
 
   let extraClassDeclaration = [{
     ::mlir::DiagnosedSilenceableFailure applyToOne(
       ::mlir::Operation *target,
       ::mlir::transform::ApplyToEachResultList &results,
       ::mlir::transform::TransformState &state);
   }];

mlir/include/mlir/Dialect/Vector/TransformOps/VectorTransformOps.td

let description = [{
 
     In other words, this makes the vector.transfer contiguous on the most minor
     dimensions and materializes the permutation_map as a vector.transpose.
   }];
 
   let assemblyFormat = "attr-dict";
 }
 
-def LowerBroadcastOp : Op<Transform_Dialect,
+def ApplyLowerBroadcastPatternsOp : Op<Transform_Dialect,
     "apply_patterns.vector.lower_broadcast",
     [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
   let description = [{
     Indicates that vector broadcast operations should be lowered to
     finer-grained vector primitives.
 
     This is usally a late step that is run after bufferization as part of the
     process of lowering to e.g. LLVM or NVVM.
   }];
 
   let assemblyFormat = "attr-dict";
 }
 
 // TODO: evolve lowering_strategy to proper enums.
-def LowerContractionOp : Op<Transform_Dialect,
+def ApplyLowerContractionPatternsOp : Op<Transform_Dialect,
     "apply_patterns.vector.lower_contraction",
     [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
   let description = [{
     Indicates that vector contraction-like operations should be lowered to
     finer-grained vector primitives.
 
     This is usually a late step that is run after bufferization as part of the
     process of lowering to e.g. LLVM or NVVM.
   }];
 
   let arguments = (ins DefaultValuedAttr<VectorContractLoweringAttr,
       "vector::VectorContractLowering::OuterProduct">:$lowering_strategy
   );
   let assemblyFormat = [{
     (`lowering_strategy` `=` $lowering_strategy^)? attr-dict
   }];
 }
 
-def LowerMasksOp : Op<Transform_Dialect,
+def ApplyLowerMasksPatternsOp : Op<Transform_Dialect,
     "apply_patterns.vector.lower_masks",
     [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
   let description = [{
     Indicates that vector.create_mask and vector.constant_mask operations
     should be lowered to finer-grained vector primitives.
 
     This is usually a late step that is run after bufferization as part of the
     process of lowering to e.g. LLVM or NVVM.
   }];
 
   let assemblyFormat = "attr-dict";
 }
 
-def LowerMaskedTransfersOp : Op<Transform_Dialect,
+def ApplyLowerMaskedTransfersPatternsOp : Op<Transform_Dialect,
     "apply_patterns.vector.lower_masked_transfers",
     [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
   let description = [{
     Indicates that masked vector.transfer and vector.gather operations should
     be lowered to finer-grained vector primitives.
 
     This is usually a late step that is run after bufferization as part of the
     process of lowering to e.g. LLVM or NVVM.
   }];
 
   let assemblyFormat = "attr-dict";
 }
 
-def MaterializeMasksOp : Op<Transform_Dialect,
+def ApplyMaterializeMasksPatternsOp : Op<Transform_Dialect,
     "apply_patterns.vector.materialize_masks",
     [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
   let description = [{
     Indicates that mask operations should be lowered to fine-grained arithemtic
     operations.
 
     This is usually the last step that is run after bufferization as part of the
     process of lowering to e.g. LLVM or NVVM.
   }];
 
   let assemblyFormat = "attr-dict";
 }
 
 // TODO: evolve lowering_strategy to proper enums.
-def LowerMultiReductionOp : Op<Transform_Dialect,
+def ApplyLowerMultiReductionPatternsOp : Op<Transform_Dialect,
     "apply_patterns.vector.lower_multi_reduction",
     [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
   let description = [{
     Indicates that vector multi_reduction-like operations should be lowered to
     finer-grained vector primitives.
 
     This is usually a late step that is run after bufferization as part of the
     process of lowering to e.g. LLVM or NVVM.
   }];
 
   let arguments = (ins DefaultValuedAttr<VectorMultiReductionLoweringAttr,
       "vector::VectorMultiReductionLowering::InnerParallel">:$lowering_strategy
   );
 
   let assemblyFormat = [{
     (`lowering_strategy` `=` $lowering_strategy^)? attr-dict
   }];
 }
 
-def LowerOuterProductOp : Op<Transform_Dialect,
+def ApplyLowerOuterProductPatternsOp : Op<Transform_Dialect,
     "apply_patterns.vector.lower_outerproduct",
     [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
   let description = [{
     Indicates that the vector outerproduct operations should be lowered to
     finer-grained vector primitives.
 
     This is usually a late step that is run after bufferization as part of the
     process of lowering to e.g. LLVM or NVVM.
   }];
 
   let assemblyFormat = "attr-dict";
 }
 
-def LowerShapeCastOp : Op<Transform_Dialect,
+def ApplyLowerGatherPatternsOp : Op<Transform_Dialect,
+    "apply_patterns.vector.lower_gather",
+    [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
+  let description = [{
+    Indicates that vector.gather operations should be lowered to
+    finer-grained vector primitives.
+  }];
+
+  let assemblyFormat = "attr-dict";
+}
+
+def ApplyLowerScanPatternsOp : Op<Transform_Dialect,
+    "apply_patterns.vector.lower_scan",
+    [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
+  let description = [{
+    Indicates that vector.scan operations should be lowered to
+    finer-grained vector primitives.
+  }];
+
+  let assemblyFormat = "attr-dict";
+}
+
+def ApplyLowerShapeCastPatternsOp : Op<Transform_Dialect,
     "apply_patterns.vector.lower_shape_cast",
     [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
   let description = [{
     Indicates that vector shape_cast operations should be lowered to
     finer-grained vector primitives.
 
     This is usually a late step that is run after bufferization as part of the
     process of lowering to e.g. LLVM or NVVM.
   }];
 
   let assemblyFormat = "attr-dict";
 }
 
-def LowerTransferOp : Op<Transform_Dialect,
+def ApplyLowerTransferPatternsOp : Op<Transform_Dialect,
     "apply_patterns.vector.lower_transfer",
     [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
   let description = [{
     Indicates that vector transfer operations should be lowered to finer-grained
     vector primitives.
 
     This is usually a late step that is run after bufferization as part of the
     process of lowering to e.g. LLVM or NVVM.
   }];
 
   let arguments = (ins DefaultValuedAttr<I64Attr, "1">:$max_transfer_rank);
 
   let assemblyFormat = [{
     (`max_transfer_rank` `=` $max_transfer_rank^)? attr-dict
   }];
 }
 
 // TODO: evolve lowering_strategy to proper enums.
-def LowerTransposeOp : Op<Transform_Dialect,
+def ApplyLowerTransposePatternsOp : Op<Transform_Dialect,
     "apply_patterns.vector.lower_transpose",
     [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
   let description = [{
     Indicates that vector transpose-like operations should be lowered to
     finer-grained vector primitives.
 
     This is usually a late step that is run after bufferization as part of the
     process of lowering to e.g. LLVM or NVVM.
oilist (
       `lowering_strategy` `=` $lowering_strategy
       | `avx2_lowering_strategy` `=` $avx2_lowering_strategy
     )
     attr-dict
   }];
 }
 
 // TODO: evolve split_transfer_strategy to proper enums.
-def SplitTransferFullPartialOp : Op<Transform_Dialect,
+def ApplySplitTransferFullPartialPatternsOp : Op<Transform_Dialect,
     "apply_patterns.vector.split_transfer_full_partial",
     [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
   let description = [{
     Indicates that vector transfer operations should be split to full and
     partial parts.
 
     This is usually a late step that is run after bufferization as part of the
     process of lowering to e.g. LLVM or NVVM.
   }];
 
   let arguments = (ins
      DefaultValuedAttr<VectorTransferSplitAttr,
        "vector::VectorTransferSplit::LinalgCopy">:$split_transfer_strategy
   );
 
   let assemblyFormat = [{
     (`split_transfer_strategy` `=` $split_transfer_strategy^)? attr-dict
   }];
 }
 
-def TransferToScfOp : Op<Transform_Dialect,
+def ApplyTransferToScfPatternsOp : Op<Transform_Dialect,
     "apply_patterns.vector.transfer_to_scf",
     [DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
   let description = [{
     Indicates that vector transfer operations should be rewritten with scf.for
     loops over finer-grained vector primitives.
 
     This is usually a late step that is run after bufferization as part of the
     process of lowering to e.g. LLVM or NVVM.

mlir/lib/Dialect/Linalg/TransformOps/DialectExtension.cpp

void init() {
     registerTransformOps<
 #define GET_OP_LIST
 #include "mlir/Dialect/Linalg/TransformOps/LinalgTransformOps.cpp.inc"
         >();
     registerTransformOps<
 #define GET_OP_LIST
 #include "mlir/Dialect/Linalg/TransformOps/LinalgMatchOps.cpp.inc"
         >();
-
-    addDialectDataInitializer<transform::PatternRegistry>(
-        [&](transform::PatternRegistry &registry) {
-          registry.registerPatterns(
-              "linalg.erase_unnecessary_inputs",
-              linalg::populateEraseUnnecessaryInputsPatterns);
-          registry.registerPatterns(
-              "linalg.fold_unit_extent_dims_via_slices",
-              linalg::populateFoldUnitExtentDimsViaSlicesPatterns);
-          registry.registerPatterns(
-              "linalg.fold_unit_extent_dims_via_reshapes",
-              linalg::populateFoldUnitExtentDimsViaReshapesPatterns);
-          registry.registerPatterns(
-              "linalg.tiling_canonicalization",
-              linalg::populateLinalgTilingCanonicalizationPatterns);
-        });
   }
 };
 } // namespace
 
 void mlir::linalg::registerTransformDialectExtension(
     DialectRegistry &registry) {
   registry.addExtensions<LinalgTransformDialectExtension>();
 }

mlir/lib/Dialect/Linalg/TransformOps/LinalgTransformOps.cpp

for (Operation *op : state.getPayloadOps(packedHandle)) {
     }
     result.push_back(op->getResult(0));
   }
 
   return DiagnosedSilenceableFailure::success();
 }
 
 //===----------------------------------------------------------------------===//
+// Apply...PatternsOp
+//===----------------------------------------------------------------------===//
+
+void transform::ApplyEraseUnnecessaryInputsPatternsOp::populatePatterns(
+    RewritePatternSet &patterns) {
+  linalg::populateEraseUnnecessaryInputsPatterns(patterns);
+}
+
+void transform::ApplyFoldUnitExtentDimsViaReshapesPatternsOp::populatePatterns(
+    RewritePatternSet &patterns) {
+  linalg::populateFoldUnitExtentDimsViaReshapesPatterns(patterns);
+}
+
+void transform::ApplyFoldUnitExtentDimsViaSlicesPatternsOp::populatePatterns(
+    RewritePatternSet &patterns) {
+  linalg::populateFoldUnitExtentDimsViaSlicesPatterns(patterns);
+}
+
+void transform::ApplyTilingCanonicalizationPatternsOp::populatePatterns(
+    RewritePatternSet &patterns) {
+  linalg::populateLinalgTilingCanonicalizationPatterns(patterns);
+}
+
+//===----------------------------------------------------------------------===//
 // BufferizeToAllocationOp
 //===----------------------------------------------------------------------===//
+
 DiagnosedSilenceableFailure
 transform::BufferizeToAllocationOp::apply(transform::TransformResults &results,
                                           transform::TransformState &state) {
   Attribute memorySpace =
       getMemorySpace().has_value() ? getMemorySpace().value() : Attribute();
   TrackingListener listener(state, *this);
   IRRewriter rewriter(getContext(), &listener);
   auto transformed = llvm::to_vector(

mlir/lib/Dialect/MemRef/TransformOps/MemRefTransformOps.cpp

 #include "mlir/Dialect/Arith/IR/Arith.h"
 #include "mlir/Dialect/MemRef/IR/MemRef.h"
 #include "mlir/Dialect/MemRef/Transforms/Passes.h"
 #include "mlir/Dialect/MemRef/Transforms/Transforms.h"
 #include "mlir/Dialect/NVGPU/IR/NVGPUDialect.h"
 #include "mlir/Dialect/SCF/IR/SCF.h"
 #include "mlir/Dialect/Transform/IR/TransformDialect.h"
 #include "mlir/Dialect/Transform/IR/TransformInterfaces.h"
-#include "mlir/Dialect/Transform/IR/TransformOps.h"
 #include "mlir/Dialect/Vector/IR/VectorOps.h"
 #include "mlir/Interfaces/LoopLikeInterface.h"
 #include "mlir/Transforms/GreedyPatternRewriteDriver.h"
 #include "llvm/Support/Debug.h"
 
 using namespace mlir;
 
 #define DEBUG_TYPE "memref-transforms"
 #define DBGS() (llvm::dbgs() << '[' << DEBUG_TYPE << "] ")
 
 //===----------------------------------------------------------------------===//
+// Apply...PatternsOp
+//===----------------------------------------------------------------------===//
+
+void transform::ApplyExpandOpsPatternsOp::populatePatterns(
+    RewritePatternSet &patterns) {
+  memref::populateExpandOpsPatterns(patterns);
+}
+
+void transform::ApplyExpandStridedMetadataPatternsOp::populatePatterns(
+    RewritePatternSet &patterns) {
+  memref::populateExpandStridedMetadataPatterns(patterns);
+}
+
+void transform::ApplyExtractAddressComputationsPatternsOp::populatePatterns(
+    RewritePatternSet &patterns) {
+  memref::populateExtractAddressComputationsPatterns(patterns);
+}
+
+void transform::ApplyFoldMemrefAliasOpsPatternsOp::populatePatterns(
+    RewritePatternSet &patterns) {
+  memref::populateFoldMemRefAliasOpPatterns(patterns);
+}
+
+void transform::ApplyResolveRankedShapedTypeResultDimsPatternsOp::
+    populatePatterns(RewritePatternSet &patterns) {
+  memref::populateResolveRankedShapedTypeResultDimsPatterns(patterns);
+}
+
+//===----------------------------------------------------------------------===//
 // MemRefMultiBufferOp
 //===----------------------------------------------------------------------===//
 
 DiagnosedSilenceableFailure transform::MemRefMultiBufferOp::apply(
     transform::TransformResults &transformResults,
     transform::TransformState &state) {
   SmallVector<Operation *> results;
   IRRewriter rewriter(getContext());
for (Operation *op : state.getPayloadOps(getTarget())) {
 
     results.push_back(*newBuffer);
   }
   transformResults.set(cast<OpResult>(getResult()), results);
   return DiagnosedSilenceableFailure::success();
 }
 
 //===----------------------------------------------------------------------===//
-// MemRefExtractAddressComputationsOp
-//===----------------------------------------------------------------------===//
-
-DiagnosedSilenceableFailure
-transform::MemRefExtractAddressComputationsOp::applyToOne(
-    Operation *target, transform::ApplyToEachResultList &results,
-    transform::TransformState &state) {
-  if (!target->hasTrait<OpTrait::IsIsolatedFromAbove>()) {
-    auto diag = this->emitOpError("requires isolated-from-above targets");
-    diag.attachNote(target->getLoc()) << "non-isolated target";
-    return DiagnosedSilenceableFailure::definiteFailure();
-  }
-
-  MLIRContext *ctx = getContext();
-  RewritePatternSet patterns(ctx);
-  memref::populateExtractAddressComputationsPatterns(patterns);
-
-  if (failed(applyPatternsAndFoldGreedily(target, std::move(patterns))))
-    return emitDefaultDefiniteFailure(target);
-
-  results.push_back(target);
-  return DiagnosedSilenceableFailure::success();
-}
-
-//===----------------------------------------------------------------------===//
 // MemRefMakeLoopIndependentOp
 //===----------------------------------------------------------------------===//
 
 DiagnosedSilenceableFailure transform::MemRefMakeLoopIndependentOp::applyToOne(
     Operation *target, transform::ApplyToEachResultList &results,
     transform::TransformState &state) {
   // Gather IVs.
   SmallVector<Value> ivs;
void init() {
     declareGeneratedDialect<memref::MemRefDialect>();
     declareGeneratedDialect<nvgpu::NVGPUDialect>();
     declareGeneratedDialect<vector::VectorDialect>();
 
     registerTransformOps<
 #define GET_OP_LIST
 #include "mlir/Dialect/MemRef/TransformOps/MemRefTransformOps.cpp.inc"
         >();
-
-    addDialectDataInitializer<transform::PatternRegistry>(
-        [&](transform::PatternRegistry &registry) {
-          registry.registerPatterns("memref.expand_ops",
-                                    memref::populateExpandOpsPatterns);
-          registry.registerPatterns("memref.fold_memref_alias_ops",
-                                    memref::populateFoldMemRefAliasOpPatterns);
-          registry.registerPatterns(
-              "memref.resolve_ranked_shaped_type_result_dims",
-              memref::populateResolveRankedShapedTypeResultDimsPatterns);
-          registry.registerPatterns(
-              "memref.expand_strided_metadata",
-              memref::populateExpandStridedMetadataPatterns);
-          registry.registerPatterns(
-              "memref.extract_address_computations",
-              memref::populateExtractAddressComputationsPatterns);
-        });
   }
 };
 } // namespace
 
 #define GET_OP_CLASSES
 #include "mlir/Dialect/MemRef/TransformOps/MemRefTransformOps.cpp.inc"
 
 void mlir::memref::registerTransformDialectExtension(
     DialectRegistry &registry) {
   registry.addExtensions<MemRefTransformDialectExtension>();
 }

mlir/lib/Dialect/SCF/TransformOps/SCFTransformOps.cpp

 #include "mlir/Dialect/Transform/IR/TransformInterfaces.h"
 #include "mlir/Dialect/Transform/IR/TransformOps.h"
 #include "mlir/Dialect/Vector/IR/VectorOps.h"
 
 using namespace mlir;
 using namespace mlir::affine;
 
 //===----------------------------------------------------------------------===//
+// Apply...PatternsOp
+//===----------------------------------------------------------------------===//
+
+void transform::ApplyForLoopCanonicalizationPatternsOp::populatePatterns(
+    RewritePatternSet &patterns) {
+  scf::populateSCFForLoopCanonicalizationPatterns(patterns);
+}
+
+//===----------------------------------------------------------------------===//
 // GetParentForOp
 //===----------------------------------------------------------------------===//
 DiagnosedSilenceableFailure
 transform::GetParentForOp::apply(transform::TransformResults &results,
                                  transform::TransformState &state) {
   SetVector<Operation *> parents;
   for (Operation *target : state.getPayloadOps(getTarget())) {
     Operation *loop, *current = target;
public:
   void init() {
     declareGeneratedDialect<affine::AffineDialect>();
     declareGeneratedDialect<func::FuncDialect>();
 
     registerTransformOps<
 #define GET_OP_LIST
 #include "mlir/Dialect/SCF/TransformOps/SCFTransformOps.cpp.inc"
         >();
-
-    addDialectDataInitializer<transform::PatternRegistry>(
-        [&](transform::PatternRegistry &registry) {
-          registry.registerPatterns(
-              "scf.for_loop_canonicalization",
-              scf::populateSCFForLoopCanonicalizationPatterns);
-        });
   }
 };
 } // namespace
 
 #define GET_OP_CLASSES
 #include "mlir/Dialect/SCF/TransformOps/SCFTransformOps.cpp.inc"
 
 void mlir::scf::registerTransformDialectExtension(DialectRegistry &registry) {
   registry.addExtensions<SCFTransformDialectExtension>();
 }

mlir/lib/Dialect/Tensor/TransformOps/TensorTransformOps.cpp

ExpandShapeOp::attachInterface<
         ReassociativeReshapeOpReplacementInterface<ExpandShapeOp>>(*ctx);
     ExtractSliceOp::attachInterface<ExtractSliceOpReplacementInterface>(*ctx);
     InsertSliceOp::attachInterface<InsertSliceOpReplacementInterface>(*ctx);
     ReshapeOp::attachInterface<ReshapeOpReplacementInterface>(*ctx);
   });
 }
 
 //===----------------------------------------------------------------------===//
+// Apply...PatternsOp
+//===----------------------------------------------------------------------===//
+
+void transform::ApplyDropRedundantInsertSliceRankExpansionPatternsOp::
+    populatePatterns(RewritePatternSet &patterns) {
+  tensor::populateDropRedundantInsertSliceRankExpansionPatterns(patterns);
+}
+
+void transform::ApplyFoldTensorEmptyPatternsOp::populatePatterns(
+    RewritePatternSet &patterns) {
+  tensor::populateFoldTensorEmptyPatterns(patterns);
+}
+
+void transform::ApplyFoldIntoPackAndUnpackPatternsOp::populatePatterns(
+    RewritePatternSet &patterns) {
+  tensor::populateFoldIntoPackAndUnpackPatterns(patterns);
+}
+
+void transform::ApplyFoldTensorSubsetOpsPatternsOp::populatePatterns(
+    RewritePatternSet &patterns) {
+  tensor::populateFoldTensorSubsetOpPatterns(patterns);
+}
+
+void transform::ApplyMergeConsecutiveInsertExtractSlicePatternsOp::
+    populatePatterns(RewritePatternSet &patterns) {
+  tensor::populateMergeConsecutiveInsertExtractSlicePatterns(patterns);
+}
+
+void transform::ApplyReassociativeReshapeFoldingPatternsOp::populatePatterns(
+    RewritePatternSet &patterns) {
+  tensor::populateReassociativeReshapeFoldingPatterns(patterns);
+}
+
+//===----------------------------------------------------------------------===//
 // MakeLoopIndependentOp
 //===----------------------------------------------------------------------===//
 
 DiagnosedSilenceableFailure transform::MakeLoopIndependentOp::applyToOne(
     Operation *target, transform::ApplyToEachResultList &results,
     transform::TransformState &state) {
   // Gather IVs.
   SmallVector<Value> ivs;
public:
   void init() {
     declareGeneratedDialect<affine::AffineDialect>();
     declareGeneratedDialect<tensor::TensorDialect>();
 
     registerTransformOps<
 #define GET_OP_LIST
 #include "mlir/Dialect/Tensor/TransformOps/TensorTransformOps.cpp.inc"
         >();
-
-    addDialectDataInitializer<transform::PatternRegistry>(
-        [&](transform::PatternRegistry &registry) {
-          registry.registerPatterns("tensor.fold_tensor_subset_ops",
-                                    tensor::populateFoldTensorSubsetOpPatterns);
-          registry.registerPatterns(
-              "tensor.merge_consecutive_insert_extract_slice",
-              tensor::populateMergeConsecutiveInsertExtractSlicePatterns);
-          registry.registerPatterns(
-              "tensor.drop_redundant_insert_slice_rank_expansion",
-              tensor::populateDropRedundantInsertSliceRankExpansionPatterns);
-          registry.registerPatterns(
-              "tensor.reassociative_reshape_folding",
-              tensor::populateReassociativeReshapeFoldingPatterns);
-          registry.registerPatterns("tensor.fold_tensor_empty",
-                                    tensor::populateFoldTensorEmptyPatterns);
-          registry.registerPatterns(
-              "tensor.fold_into_pack_and_unpack",
-              tensor::populateFoldIntoPackAndUnpackPatterns);
-        });
   }
 };
 } // namespace
 
 #define GET_OP_CLASSES
 #include "mlir/Dialect/Tensor/TransformOps/TensorTransformOps.cpp.inc"
 
 void mlir::tensor::registerTransformDialectExtension(
     DialectRegistry &registry) {
   registry.addExtensions<TensorTransformDialectExtension>();
 }

mlir/lib/Dialect/Transform/IR/TransformOps.cpp

 #define DBGS() (llvm::dbgs() << "[" DEBUG_TYPE "] ")
 
 #define DEBUG_TYPE_MATCHER "transform-matcher"
 #define DBGS_MATCHER() (llvm::dbgs() << "[" DEBUG_TYPE_MATCHER "] ")
 #define DEBUG_MATCHER(x) DEBUG_WITH_TYPE(DEBUG_TYPE_MATCHER, x)
 
 using namespace mlir;
 
-MLIR_DEFINE_EXPLICIT_TYPE_ID(mlir::transform::PatternRegistry)
-
 static ParseResult parseSequenceOpOperands(
     OpAsmParser &parser, std::optional<OpAsmParser::UnresolvedOperand> &root,
     Type &rootType,
     SmallVectorImpl<OpAsmParser::UnresolvedOperand> &extraBindings,
     SmallVectorImpl<Type> &extraBindingTypes);
 static void printSequenceOpOperands(OpAsmPrinter &printer, Operation *op,
                                     Value root, Type rootType,
                                     ValueRange extraBindings,
void transform::ErrorCheckingTrackingListener::notifyPayloadReplacementNotFound(
   for (auto &&[index, value] : llvm::enumerate(values))
     status.attachNote(value.getLoc())
         << "[" << errorCounter << "] replacement value " << index;
 
   ++errorCounter;
 }
 
 //===----------------------------------------------------------------------===//
-// PatternRegistry
-//===----------------------------------------------------------------------===//
-
-void transform::PatternRegistry::registerPatterns(StringRef identifier,
-                                                  PopulatePatternsFn &&fn) {
-  StringAttr attr = builder.getStringAttr(identifier);
-  assert(!patterns.contains(attr) && "patterns identifier is already in use");
-  patterns.try_emplace(attr, std::move(fn));
-}
-
-void transform::PatternRegistry::registerPatterns(
-    StringRef identifier, PopulatePatternsWithBenefitFn &&fn) {
-  StringAttr attr = builder.getStringAttr(identifier);
-  assert(!patterns.contains(attr) && "patterns identifier is already in use");
-  patterns.try_emplace(attr, [f = std::move(fn)](RewritePatternSet &patternSet) {
-    f(patternSet, /*benefit=*/1);
-  });
-}
-
-void transform::PatternRegistry::populatePatterns(
-    StringAttr identifier, RewritePatternSet &patternSet) const {
-  auto it = patterns.find(identifier);
-  assert(it != patterns.end() && "patterns not registered in registry");
-  it->second(patternSet);
-}
-
-bool transform::PatternRegistry::hasPatterns(StringAttr identifier) const {
-  return patterns.contains(identifier);
-}
-
-//===----------------------------------------------------------------------===//
 // AlternativesOp
 //===----------------------------------------------------------------------===//
 
 OperandRange transform::AlternativesOp::getSuccessorEntryOperands(
     std::optional<unsigned> index) {
   if (index && getOperation()->getNumOperands() == 1)
     return getOperation()->getOperands();
   return OperandRange(getOperation()->operand_end(),

 
 DiagnosedSilenceableFailure
 transform::ApplyPatternsOp::applyToOne(Operation *target,
                                        ApplyToEachResultList &results,
                                        transform::TransformState &state) {
   // Gather all specified patterns.
   MLIRContext *ctx = target->getContext();
   RewritePatternSet patterns(ctx);
-  const auto &registry = getContext()
-                             ->getLoadedDialect<transform::TransformDialect>()
-                             ->getExtraData<transform::PatternRegistry>();
-  for (Attribute attr : getPatterns())
-    registry.populatePatterns(attr.cast<StringAttr>(), patterns);
   if (!getRegion().empty()) {
     for (Operation &op : getRegion().front()) {
       cast<transform::PatternDescriptorOpInterface>(&op).populatePatterns(
           patterns);
     }
   }
 
   // Configure the GreedyPatternRewriteDriver.
if (getFailOnPayloadReplacementNotFound())
       return status;
     (void)status.silence();
   }
 
   return DiagnosedSilenceableFailure::success();
 }
 
 LogicalResult transform::ApplyPatternsOp::verify() {
-  const auto &registry = getContext()
-                             ->getLoadedDialect<transform::TransformDialect>()
-                             ->getExtraData<transform::PatternRegistry>();
-  for (Attribute attr : getPatterns()) {
-    auto strAttr = attr.dyn_cast<StringAttr>();
-    if (!strAttr)
-      return emitOpError() << "expected " << getPatternsAttrName()
-                           << " to be an array of strings";
-    if (!registry.hasPatterns(strAttr))
-      return emitOpError() << "patterns not registered: " << strAttr.strref();
-  }
   if (!getRegion().empty()) {
     for (Operation &op : getRegion().front()) {
       if (!isa<transform::PatternDescriptorOpInterface>(&op)) {
         InFlightDiagnostic diag = emitOpError()
                                   << "expected children ops to implement "
                                      "PatternDescriptorOpInterface";
         diag.attachNote(op.getLoc()) << "op without interface";
         return diag;

mlir/lib/Dialect/Vector/TransformOps/VectorTransformOps.cpp

 #include "mlir/Dialect/X86Vector/Transforms.h"
 #include "mlir/Transforms/GreedyPatternRewriteDriver.h"
 
 using namespace mlir;
 using namespace mlir::vector;
 using namespace mlir::transform;
 
 //===----------------------------------------------------------------------===//
-// ApplyRankReducingSubviewPatternsOp
+// Apply...PatternsOp
 //===----------------------------------------------------------------------===//
 
 void transform::ApplyRankReducingSubviewPatternsOp::populatePatterns(
     RewritePatternSet &patterns) {
   vector::populateVectorTransferDropUnitDimsPatterns(patterns);
 }
 
-//===----------------------------------------------------------------------===//
-// ApplyTransferPermutationPatternsOp
-//===----------------------------------------------------------------------===//
-
 void transform::ApplyTransferPermutationPatternsOp::populatePatterns(
     RewritePatternSet &patterns) {
   vector::populateVectorTransferPermutationMapLoweringPatterns(patterns);
 }
 
-//===----------------------------------------------------------------------===//
-// LowerBroadcastOp
-//===----------------------------------------------------------------------===//
-
-void transform::LowerBroadcastOp::populatePatterns(
+void transform::ApplyLowerBroadcastPatternsOp::populatePatterns(
     RewritePatternSet &patterns) {
   populateVectorBroadcastLoweringPatterns(patterns);
 }
 
-//===----------------------------------------------------------------------===//
-// LowerContractionOp
-//===----------------------------------------------------------------------===//
-
-void transform::LowerContractionOp::populatePatterns(
+void transform::ApplyLowerContractionPatternsOp::populatePatterns(
     RewritePatternSet &patterns) {
   vector::VectorTransformsOptions vectorTransformOptions;
   vectorTransformOptions.setVectorTransformsOptions(getLoweringStrategy());
   populateVectorContractLoweringPatterns(patterns, vectorTransformOptions,
                                          /*benefit=*/1,
                                          /*disableOuterProductLowering=*/true);
 }
 
-//===----------------------------------------------------------------------===//
-// LowerMasksOp
-//===----------------------------------------------------------------------===//
-
-void transform::LowerMasksOp::populatePatterns(RewritePatternSet &patterns) {
+void transform::ApplyLowerMasksPatternsOp::populatePatterns(
+    RewritePatternSet &patterns) {
   populateVectorMaskOpLoweringPatterns(patterns);
 }
 
-//===----------------------------------------------------------------------===//
-// LowerMaskedTransfersOp
-//===----------------------------------------------------------------------===//
-
-void transform::LowerMaskedTransfersOp::populatePatterns(
+void transform::ApplyLowerMaskedTransfersPatternsOp::populatePatterns(
     RewritePatternSet &patterns) {
   populateVectorMaskLoweringPatternsForSideEffectingOps(patterns);
 }
 
-//===----------------------------------------------------------------------===//
-// MaterializeMasksOp
-//===----------------------------------------------------------------------===//
-
-void transform::MaterializeMasksOp::populatePatterns(RewritePatternSet &patterns) {
+void transform::ApplyMaterializeMasksPatternsOp::populatePatterns(
+    RewritePatternSet &patterns) {
   populateVectorMaskMaterializationPatterns(patterns,
                                             /*force32BitVectorIndices=*/false);
 }
 
-//===----------------------------------------------------------------------===//
-// LowerMultiReductionOp
-//===----------------------------------------------------------------------===//
-
-void transform::LowerMultiReductionOp::populatePatterns(
+void transform::ApplyLowerMultiReductionPatternsOp::populatePatterns(
     RewritePatternSet &patterns) {
   vector::VectorTransformsOptions vectorTransformOptions;
   vectorTransformOptions.setVectorMultiReductionLowering(getLoweringStrategy());
   vector::populateVectorMultiReductionLoweringPatterns(
       patterns, vectorTransformOptions.vectorMultiReductionLowering);
 }
 
-//===----------------------------------------------------------------------===//
-// LowerOuterProductOp
-//===----------------------------------------------------------------------===//
-
-void transform::LowerOuterProductOp::populatePatterns(
+void transform::ApplyLowerOuterProductPatternsOp::populatePatterns(
     RewritePatternSet &patterns) {
   populateVectorOuterProductLoweringPatterns(patterns);
 }
 
-//===----------------------------------------------------------------------===//
-// LowerShapeCastOp
-//===----------------------------------------------------------------------===//
+void transform::ApplyLowerGatherPatternsOp::populatePatterns(
+    RewritePatternSet &patterns) {
+  vector::populateVectorGatherLoweringPatterns(patterns);
+}
 
-void transform::LowerShapeCastOp::populatePatterns(
+void transform::ApplyLowerScanPatternsOp::populatePatterns(
     RewritePatternSet &patterns) {
-  vector::populateVectorShapeCastLoweringPatterns(patterns);
+  vector::populateVectorScanLoweringPatterns(patterns);
 }
 
-//===----------------------------------------------------------------------===//
-// LowerTransferOp
-//===----------------------------------------------------------------------===//
+void transform::ApplyLowerShapeCastPatternsOp::populatePatterns(
+    RewritePatternSet &patterns) {
+  vector::populateVectorShapeCastLoweringPatterns(patterns);
+}
 
-void transform::LowerTransferOp::populatePatterns(RewritePatternSet &patterns) {
+void transform::ApplyLowerTransferPatternsOp::populatePatterns(
+    RewritePatternSet &patterns) {
   vector::populateVectorTransferLoweringPatterns(patterns,
                                                  getMaxTransferRank());
 }
 
-//===----------------------------------------------------------------------===//
-// LowerTransposeOp
-//===----------------------------------------------------------------------===//
-
-void transform::LowerTransposeOp::populatePatterns(
+void transform::ApplyLowerTransposePatternsOp::populatePatterns(
     RewritePatternSet &patterns) {
   vector::populateVectorTransposeLoweringPatterns(
       patterns, vector::VectorTransformsOptions().setVectorTransposeLowering(
                     getLoweringStrategy()));
   if (getAvx2LoweringStrategy()) {
     auto avx2LoweringOptions =
         x86vector::avx2::LoweringOptions().setTransposeOptions(
             x86vector::avx2::TransposeLoweringOptions()
                 .lower4x8xf32(true)
                 .lower8x8xf32(true));
     x86vector::avx2::populateSpecializedTransposeLoweringPatterns(
         patterns, avx2LoweringOptions, /*benefit=*/10);
   }
 }
 
-//===----------------------------------------------------------------------===//
-// SplitTransferFullPartialOp
-//===----------------------------------------------------------------------===//
-
-void transform::SplitTransferFullPartialOp::populatePatterns(
+void transform::ApplySplitTransferFullPartialPatternsOp::populatePatterns(
     RewritePatternSet &patterns) {
   vector::VectorTransformsOptions vectorTransformOptions;
   vectorTransformOptions.setVectorTransferSplit(getSplitTransferStrategy());
   populateVectorTransferFullPartialPatterns(patterns, vectorTransformOptions);
 }
 
-//===----------------------------------------------------------------------===//
-// TransferToScfOp
-//===----------------------------------------------------------------------===//
-
-void transform::TransferToScfOp::populatePatterns(RewritePatternSet &patterns) {
+void transform::ApplyTransferToScfPatternsOp::populatePatterns(
+    RewritePatternSet &patterns) {
   VectorTransferToSCFOptions vectorTransferToSCFOptions =
       VectorTransferToSCFOptions()
           .enableFullUnroll(getFullUnroll())
           .setTargetRank(getMaxTransferRank());
   populateVectorToSCFConversionPatterns(patterns, vectorTransferToSCFOptions);
 }
 
 //===----------------------------------------------------------------------===//

 public:
   VectorTransformDialectExtension() {
     declareGeneratedDialect<vector::VectorDialect>();
     declareGeneratedDialect<LLVM::LLVMDialect>();
     registerTransformOps<
 #define GET_OP_LIST
 #include "mlir/Dialect/Vector/TransformOps/VectorTransformOps.cpp.inc"
         >();
-
-    addDialectDataInitializer<transform::PatternRegistry>(
-        [&](transform::PatternRegistry &registry) {
-          registry.registerPatterns("vector.outer_product_lowering",
-                                    populateVectorOuterProductLoweringPatterns);
-          registry.registerPatterns("vector.broadcast_lowering",
-                                    populateVectorBroadcastLoweringPatterns);
-          registry.registerPatterns("vector.mask_op_lowering",
-                                    populateVectorMaskOpLoweringPatterns);
-          registry.registerPatterns("vector.shape_cast_lowering",
-                                    populateVectorShapeCastLoweringPatterns);
-          registry.registerPatterns(
-              "vector.transfer_lowering",
-              [&](RewritePatternSet &set, PatternBenefit benefit) {
-                return populateVectorTransferLoweringPatterns(
-                    set, /*maxTransferRank=*/std::nullopt, benefit);
-              });
-          registry.registerPatterns(
-              "vector.transfer_permutation_map_lowering",
-              populateVectorTransferPermutationMapLoweringPatterns);
-          registry.registerPatterns("vector.scan_lowering",
-                                    populateVectorScanLoweringPatterns);
-          registry.registerPatterns("vector.vector_gather_lowering",
-                                    populateVectorGatherLoweringPatterns);
-          registry.registerPatterns(
-              "vector.mask_lowering_for_side_effecting_ops",
-              populateVectorMaskLoweringPatternsForSideEffectingOps);
-        });
   }
 };
 } // namespace
 
 #define GET_OP_CLASSES
 #include "mlir/Dialect/Vector/TransformOps/VectorTransformOps.cpp.inc"
 
 void mlir::vector::registerTransformDialectExtension(
     DialectRegistry &registry) {
   registry.addExtensions<VectorTransformDialectExtension>();
 }

mlir/test/Dialect/LLVM/transform-e2e.mlir

%b = transform.bufferization.one_shot_bufferize layout{IdentityLayoutMap}
       %module_op {bufferize_function_boundaries = true}
       : (!transform.any_op) -> !transform.any_op
 
   %f = transform.structured.match ops{["func.func"]} in %b
     : (!transform.any_op) -> !transform.any_op
 
   // TODO: group these lower-level controls into various properly named vector
   // lowering TD macros.
-  transform.apply_patterns [] to %f {
+  transform.apply_patterns to %f {
     transform.apply_patterns.vector.lower_contraction lowering_strategy = "outerproduct"
   } : !transform.any_op
 
-  transform.apply_patterns [] to %f {
+  transform.apply_patterns to %f {
     transform.apply_patterns.vector.apply_transfer_permutation_patterns
   } : !transform.any_op
 
-  transform.apply_patterns [] to %f {
+  transform.apply_patterns to %f {
     transform.apply_patterns.vector.lower_multi_reduction lowering_strategy = "innerparallel"
   } : !transform.any_op
 
-  transform.apply_patterns [] to %f {
+  transform.apply_patterns to %f {
     transform.apply_patterns.vector.split_transfer_full_partial split_transfer_strategy = "linalg-copy"
   } : !transform.any_op
 
-  transform.apply_patterns [] to %f {
+  transform.apply_patterns to %f {
     transform.apply_patterns.vector.transfer_to_scf max_transfer_rank = 1 full_unroll = true
   } : !transform.any_op
 
-  transform.apply_patterns [] to %f {
+  transform.apply_patterns to %f {
     transform.apply_patterns.vector.lower_transfer max_transfer_rank = 1
   } : !transform.any_op
 
-  transform.apply_patterns [] to %f {
+  transform.apply_patterns to %f {
     transform.apply_patterns.vector.lower_shape_cast
   } : !transform.any_op
 
-  transform.apply_patterns [] to %f {
+  transform.apply_patterns to %f {
     transform.apply_patterns.vector.lower_transpose lowering_strategy = "shuffle_1d"
   } : !transform.any_op
 }

mlir/test/Dialect/MemRef/extract-address-computations.mlir

func.func @test_load(%base : memref<2x16x16xf32>, %offset : index) -> f32 {
   %c8 = arith.constant 8 : index
   %loaded_val = memref.load %base[%offset, %c0, %c8] : memref<2x16x16xf32>
   return %loaded_val : f32
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%arg1: !transform.any_op):
   %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!transform.any_op) -> !transform.any_op
-  %1 = transform.memref.extract_address_computations %0 : (!transform.any_op) -> !transform.any_op
+  transform.apply_patterns to %0 {
+    transform.apply_patterns.memref.extract_address_computations
+  } : !transform.any_op
   // Verify that the returned handle is usable.
-  transform.test_print_remark_at_operand %1, "transformed" : !transform.any_op
+  transform.test_print_remark_at_operand %0, "transformed" : !transform.any_op
 }
 
 // -----
 
 // Same as previous @test_load but with the nontemporal flag.
 
 // CHECK-LABEL: @test_load_nontemporal(
 // CHECK-SAME: %[[BASE:[^:]*]]: memref{{[^,]*}},
 // CHECK-SAME: %[[DYN_OFFSET:.*]]: index)
 // CHECK-DAG: %[[C0:.*]] = arith.constant 0 : index
 // CHECK-DAG: %[[SUBVIEW:.*]] = memref.subview %[[BASE]][%[[DYN_OFFSET]], 0, 8] [1, 1, 1] [1, 1, 1] : memref<2x16x16xf32> to memref<1x1x1xf32, strided<[256, 16, 1], offset: ?>>
 // CHECK: %[[LOADED_VAL:.*]] = memref.load %[[SUBVIEW]][%[[C0]], %[[C0]], %[[C0]]] {nontemporal = true} : memref<1x1x1xf32, strided<[256, 16, 1], offset: ?>>
 // CHECK: return %[[LOADED_VAL]] : f32
 func.func @test_load_nontemporal(%base : memref<2x16x16xf32>, %offset : index) -> f32 {
   %c0 = arith.constant 0 : index
   %c8 = arith.constant 8 : index
   %loaded_val = memref.load %base[%offset, %c0, %c8] {nontemporal = true } : memref<2x16x16xf32>
   return %loaded_val : f32
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%arg1: !transform.any_op):
   %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!transform.any_op) -> !transform.any_op
-  %1 = transform.memref.extract_address_computations %0 : (!transform.any_op) -> !transform.any_op
-}
+  transform.apply_patterns to %0 {
+    transform.apply_patterns.memref.extract_address_computations
+  } : !transform.any_op}
 
 // -----
 
 // Simple test: check that we extract the address computation of a store into
 // a dedicated subview.
 // The resulting store will use the address from the subview and have only
 // indices set to zero.
 
func.func @test_store(%base : memref<2x16x16xf32>, %offset : index) -> () {
   %c8 = arith.constant 8 : index
   memref.store %cf0, %base[%offset, %c0, %c8] : memref<2x16x16xf32>
   return
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%arg1: !transform.any_op):
   %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!transform.any_op) -> !transform.any_op
-  %1 = transform.memref.extract_address_computations %0 : (!transform.any_op) -> !transform.any_op
-}
+  transform.apply_patterns to %0 {
+    transform.apply_patterns.memref.extract_address_computations
+  } : !transform.any_op}
 
 // -----
 
 // Same as @test_store but check that the nontemporal flag is preserved.
 
 // CHECK-LABEL: @test_store_nontemporal(
 // CHECK-SAME: %[[BASE:[^:]*]]: memref{{[^,]*}},
 // CHECK-SAME: %[[DYN_OFFSET:.*]]: index)
 // CHECK-DAG: %[[CF0:.*]] = arith.constant 0.0{{0*e\+00}} : f32
 // CHECK-DAG: %[[C0:.*]] = arith.constant 0 : index
 // CHECK-DAG: %[[SUBVIEW:.*]] = memref.subview %[[BASE]][%[[DYN_OFFSET]], 0, 8] [1, 1, 1] [1, 1, 1] : memref<2x16x16xf32> to memref<1x1x1xf32, strided<[256, 16, 1], offset: ?>>
 // CHECK: memref.store %[[CF0]], %[[SUBVIEW]][%[[C0]], %[[C0]], %[[C0]]] {nontemporal = true} : memref<1x1x1xf32, strided<[256, 16, 1], offset: ?>>
 // CHECK: return
 func.func @test_store_nontemporal(%base : memref<2x16x16xf32>, %offset : index) -> () {
   %cf0 = arith.constant 0.0 : f32
   %c0 = arith.constant 0 : index
   %c8 = arith.constant 8 : index
   memref.store %cf0, %base[%offset, %c0, %c8] { nontemporal = true } : memref<2x16x16xf32>
   return
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%arg1: !transform.any_op):
   %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!transform.any_op) -> !transform.any_op
-  %1 = transform.memref.extract_address_computations %0 : (!transform.any_op) -> !transform.any_op
-}
+  transform.apply_patterns to %0 {
+    transform.apply_patterns.memref.extract_address_computations
+  } : !transform.any_op}
 
 // -----
 // For this test, we made the source memref fully dynamic.
 // The gist of the check remains the same as the simple test:
 // The address computation is extracted into its own subview.
 // CHECK-LABEL: @testWithLoop(
 // CHECK-SAME: %[[BASE:[^:]*]]: memref
 // CHECK:  %[[SUM_ALL:.*]] = arith.constant 0.0{{0*e\+00}} : f32
%sum_res2 = scf.for %iv2 = %c0 to %upper_bound2 step %c1 iter_args(%sum_iter2 = %sum_all) -> (f32) {
     scf.yield %sum_res1 : f32
   }
   return %sum_res2 : f32
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%arg1: !transform.any_op):
   %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!transform.any_op) -> !transform.any_op
-  %1 = transform.memref.extract_address_computations %0 : (!transform.any_op) -> !transform.any_op
-}
+  transform.apply_patterns to %0 {
+    transform.apply_patterns.memref.extract_address_computations
+  } : !transform.any_op}
 
 // -----
 
 // Simple test: check that we extract the address computation of a ldmatrix into
 // a dedicated subview.
 // The resulting ldmatrix will loaded from with subview and have only indices set
 // to zero.
 // Also the sizes of the view are adjusted to `original size - offset`.
%loaded_val = nvgpu.ldmatrix
     {numTiles = 4 : i32, transpose = false}
       : memref<4x32x32xf16, 3> -> vector<4x2xf16>
   return %loaded_val : vector<4x2xf16>
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%arg1: !transform.any_op):
   %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!transform.any_op) -> !transform.any_op
-  %1 = transform.memref.extract_address_computations %0 : (!transform.any_op) -> !transform.any_op
-}
+  transform.apply_patterns to %0 {
+    transform.apply_patterns.memref.extract_address_computations
+  } : !transform.any_op}
 
 // -----
 
 // Same as test_ldmatrix but with fully dynamic memref.
 
 // CHECK-DAG: #[[$A_MINUS_B_MAP:.*]] = affine_map<()[s0, s1] -> (s0 - s1)>
 // CHECK-LABEL: @test_ldmatrix(
 // CHECK-SAME: %[[BASE:[^:]*]]: memref<{{[^,]*}}, 3>,
%loaded_val = nvgpu.ldmatrix
     {numTiles = 4 : i32, transpose = false}
       : memref<?x?x?xf16, 3> -> vector<4x2xf16>
   return %loaded_val : vector<4x2xf16>
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%arg1: !transform.any_op):
   %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!transform.any_op) -> !transform.any_op
-  %1 = transform.memref.extract_address_computations %0 : (!transform.any_op) -> !transform.any_op
-}
+  transform.apply_patterns to %0 {
+    transform.apply_patterns.memref.extract_address_computations
+  } : !transform.any_op}
 
 // -----
 
 // Simple test for vector.transfer_read with fully dynamic memref.
 // We also set a permutation map to make sure it is properly preserved.
 
 // CHECK-DAG: #[[$A_MINUS_B_MAP:.*]] = affine_map<()[s0, s1] -> (s0 - s1)>
 // CHECK-DAG: #[[$PERMUTATION_MAP:.*]] = affine_map<(d0, d1, d2) -> (d2, d0)>
func.func @test_transfer_read_op(%base : memref<?x?x?xf16>,
   %cf0 = arith.constant 0.0 : f16
   %loaded_val = vector.transfer_read %base[%offset0, %offset1, %offset2], %cf0 { permutation_map = affine_map<(d0,d1,d2) -> (d2,d0)> } : memref<?x?x?xf16>, vector<4x2xf16>
   return %loaded_val : vector<4x2xf16>
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%arg1: !transform.any_op):
   %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!transform.any_op) -> !transform.any_op
-  %1 = transform.memref.extract_address_computations %0 : (!transform.any_op) -> !transform.any_op
-}
+  transform.apply_patterns to %0 {
+    transform.apply_patterns.memref.extract_address_computations
+  } : !transform.any_op}
 
 // -----
 
 // Same as test_transfer_read_op but with tensors.
 // Right now this rewrite is not supported but we still shouldn't choke on it.
 
 // CHECK: #[[$PERMUTATION_MAP:.*]] = affine_map<(d0, d1, d2) -> (d2, d0)>
 // CHECK-LABEL: @test_transfer_read_op_with_tensor(
func.func @test_transfer_read_op_with_tensor(%base : tensor<?x?x?xf16>,
   %cf0 = arith.constant 0.0 : f16
   %loaded_val = vector.transfer_read %base[%offset0, %offset1, %offset2], %cf0 { permutation_map = affine_map<(d0,d1,d2) -> (d2,d0)> } : tensor<?x?x?xf16>, vector<4x2xf16>
   return %loaded_val : vector<4x2xf16>
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%arg1: !transform.any_op):
   %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!transform.any_op) -> !transform.any_op
-  %1 = transform.memref.extract_address_computations %0 : (!transform.any_op) -> !transform.any_op
-}
+  transform.apply_patterns to %0 {
+    transform.apply_patterns.memref.extract_address_computations
+  } : !transform.any_op}
 
 // -----
 
 // Simple test for vector.transfer_write with fully dynamic memref.
 // We also set a permutation map to make sure it is properly preserved.
 
 // CHECK-DAG: #[[$A_MINUS_B_MAP:.*]] = affine_map<()[s0, s1] -> (s0 - s1)>
 // CHECK-DAG: #[[$PERMUTATION_MAP:.*]] = affine_map<(d0, d1, d2) -> (d2, d0)>
func.func @test_transfer_write_op(%base : memref<?x?x?xf16>,
   %vcf0 = arith.constant dense<0.000000e+00> : vector<4x2xf16>
   vector.transfer_write %vcf0, %base[%offset0, %offset1, %offset2] { permutation_map = affine_map<(d0,d1,d2) -> (d2,d0)> } : vector<4x2xf16>, memref<?x?x?xf16>
   return
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%arg1: !transform.any_op):
   %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!transform.any_op) -> !transform.any_op
-  %1 = transform.memref.extract_address_computations %0 : (!transform.any_op) -> !transform.any_op
-}
+  transform.apply_patterns to %0 {
+    transform.apply_patterns.memref.extract_address_computations
+  } : !transform.any_op}
 
 // -----
 
 // Check that the strides of the original memref are kept.
 // Moreover even with non-1 strides the subview should still issue [1,...]
 // strides, since this is a multiplication factor.
 
 // CHECK-DAG: #[[$A_MINUS_B_MAP:.*]] = affine_map<()[s0, s1] -> (s0 - s1)>
func.func @test_transfer_write_op_with_strides(%base : memref<?x?x?xf16, strided<[329, 26, 12], offset: ?>>,
   %vcf0 = arith.constant dense<0.000000e+00> : vector<4x2xf16>
   vector.transfer_write %vcf0, %base[%offset0, %offset1, %offset2] { permutation_map = affine_map<(d0,d1,d2) -> (d2,d0)> } : vector<4x2xf16>, memref<?x?x?xf16, strided<[329, 26, 12], offset: ?>>
   return
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%arg1: !transform.any_op):
   %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!transform.any_op) -> !transform.any_op
-  %1 = transform.memref.extract_address_computations %0 : (!transform.any_op) -> !transform.any_op
-}
+  transform.apply_patterns to %0 {
+    transform.apply_patterns.memref.extract_address_computations
+  } : !transform.any_op}
 // -----
 
 // Same as test_transfer_write_op but with tensors.
 // Right now this rewrite is not supported but we still shouldn't choke on it.
 
 // CHECK: #[[$PERMUTATION_MAP:.*]] = affine_map<(d0, d1, d2) -> (d2, d0)>
 // CHECK-LABEL: @test_transfer_write_op_with_tensor(
 // CHECK-SAME: %[[BASE:[^:]*]]: tensor<{{[^,]*}}>,
 // CHECK-SAME: %[[DYN_OFFSET0:[^:]*]]: index,
 // CHECK-SAME: %[[DYN_OFFSET1:[^:]*]]: index,
 // CHECK-SAME: %[[DYN_OFFSET2:[^:]*]]: index)
 // CHECK-DAG: %[[VCF0:.*]] = arith.constant dense<0.0{{0*e\+00}}> : vector<4x2xf16>
 // CHECK: %[[RES:.*]] = vector.transfer_write %[[VCF0]], %[[BASE]][%[[DYN_OFFSET0]], %[[DYN_OFFSET1]], %[[DYN_OFFSET2]]] {permutation_map = #[[$PERMUTATION_MAP]]} : vector<4x2xf16>, tensor<?x?x?xf16>
 // CHECK: return %[[RES]] : tensor<?x?x?xf16>
 func.func @test_transfer_write_op_with_tensor(%base : tensor<?x?x?xf16>,
     %offset0 : index, %offset1: index, %offset2: index) -> tensor<?x?x?xf16> {
   %vcf0 = arith.constant dense<0.000000e+00> : vector<4x2xf16>
   %res = vector.transfer_write %vcf0, %base[%offset0, %offset1, %offset2] { permutation_map = affine_map<(d0,d1,d2) -> (d2,d0)> } : vector<4x2xf16>, tensor<?x?x?xf16>
   return %res : tensor<?x?x?xf16>
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%arg1: !transform.any_op):
   %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!transform.any_op) -> !transform.any_op
-  %1 = transform.memref.extract_address_computations %0 : (!transform.any_op) -> !transform.any_op
-}
+  transform.apply_patterns to %0 {
+    transform.apply_patterns.memref.extract_address_computations
+  } : !transform.any_op}

mlir/test/Dialect/Transform/ops-invalid.mlir

^bb0(%arg0: !transform.any_op):
   // expected-error @below {{expected "patterns" to be an array of strings}}
   transform.apply_patterns [3, 9] to %arg0 {} : !transform.any_op
 }
 
 // -----
 transform.sequence failures(propagate) {
 ^bb0(%arg0: !transform.any_op):
   // expected-error @below {{expected children ops to implement PatternDescriptorOpInterface}}
-  transform.apply_patterns [] to %arg0 {
+  transform.apply_patterns to %arg0 {
     // expected-note @below {{op without interface}}
     transform.named_sequence @foo()
   } : !transform.any_op
 }

mlir/test/Dialect/Transform/test-pattern-application.mlir

func.func @update_tracked_op_mapping() {
   }) : () -> ()
   return
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%arg1: !transform.any_op):
   %0 = transform.structured.match ops{["test.container"]} in %arg1 : (!transform.any_op) -> !transform.any_op
   %1 = transform.structured.match ops{["test.foo"]} in %arg1 : (!transform.any_op) -> !transform.any_op
-  transform.apply_patterns ["transform.test"] to %0 {} : !transform.any_op
-  // Add an attribute to %1, which is now mapped to a new op.
-  transform.annotate %1 "annotated" : !transform.any_op
-}
-
-// -----
-
-// CHECK-LABEL: func @update_tracked_op_mapping_region()
-//       CHECK:   "test.container"() ({
-//       CHECK:     %0 = "test.foo"() {annotated} : () -> i32
-//       CHECK:   }) : () -> ()
-func.func @update_tracked_op_mapping_region() {
-  "test.container"() ({
-    %0 = "test.foo"() {replace_with_new_op = "test.foo"} : () -> (i32)
-  }) : () -> ()
-  return
-}
-
-transform.sequence failures(propagate) {
-^bb1(%arg1: !transform.any_op):
-  %0 = transform.structured.match ops{["test.container"]} in %arg1 : (!transform.any_op) -> !transform.any_op
-  %1 = transform.structured.match ops{["test.foo"]} in %arg1 : (!transform.any_op) -> !transform.any_op
-  transform.apply_patterns [] to %0 {
+  transform.apply_patterns to %0 {
     transform.apply_patterns.transform.test_patterns
   } : !transform.any_op
   // Add an attribute to %1, which is now mapped to a new op.
   transform.annotate %1 "annotated" : !transform.any_op
 }
 
 // -----
 
 func.func @replacement_op_not_found() {
   "test.container"() ({
     // expected-note @below {{[0] replaced op}}
     // expected-note @below {{[0] replacement value 0}}
     %0 = "test.foo"() {replace_with_new_op = "test.bar"} : () -> (i32)
   }) : () -> ()
   return
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%arg1: !transform.any_op):
   %0 = transform.structured.match ops{["test.container"]} in %arg1 : (!transform.any_op) -> !transform.any_op
   %1 = transform.structured.match ops{["test.foo"]} in %arg1 : (!transform.any_op) -> !transform.any_op
   // expected-error @below {{tracking listener failed to find replacement op}}
-  transform.apply_patterns ["transform.test"] to %0 {} : !transform.any_op
+  transform.apply_patterns to %0 {
+    transform.apply_patterns.transform.test_patterns
+  } : !transform.any_op
   // %1 must be used in some way. If no replacement payload op could be found,
   // an error is thrown only if the handle is not dead.
   transform.annotate %1 "annotated" : !transform.any_op
 }
 
 // -----
 
 // CHECK-LABEL: func @replacement_op_for_dead_handle_not_found()
 //       CHECK:   "test.container"() ({
 //       CHECK:     %0 = "test.bar"() : () -> i32
 //       CHECK:   }) : () -> ()
 func.func @replacement_op_for_dead_handle_not_found() {
   "test.container"() ({
     %0 = "test.foo"() {replace_with_new_op = "test.bar"} : () -> (i32)
   }) : () -> ()
   return
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%arg1: !transform.any_op):
   %0 = transform.structured.match ops{["test.container"]} in %arg1 : (!transform.any_op) -> !transform.any_op
   %1 = transform.structured.match ops{["test.foo"]} in %arg1 : (!transform.any_op) -> !transform.any_op
   // No error because %1 is dead.
-  transform.apply_patterns ["transform.test"] to %0 {} : !transform.any_op
+  transform.apply_patterns to %0 {
+    transform.apply_patterns.transform.test_patterns
+  } : !transform.any_op
 }
 
 // -----
 
 // CHECK-LABEL: func @replacement_op_not_found_silenced()
 //       CHECK:   "test.container"() ({
 //       CHECK:     %0 = "test.bar"() : () -> i32
 //       CHECK:   }) : () -> ()
 func.func @replacement_op_not_found_silenced() {
   "test.container"() ({
     %0 = "test.foo"() {replace_with_new_op = "test.bar"} : () -> (i32)
   }) : () -> ()
   return
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%arg1: !transform.any_op):
   %0 = transform.structured.match ops{["test.container"]} in %arg1 : (!transform.any_op) -> !transform.any_op
   %1 = transform.structured.match ops{["test.foo"]} in %arg1 : (!transform.any_op) -> !transform.any_op
-  transform.apply_patterns ["transform.test"] to %0 {} {fail_on_payload_replacement_not_found = false}: !transform.any_op
+  transform.apply_patterns to %0 {
+    transform.apply_patterns.transform.test_patterns
+  } {fail_on_payload_replacement_not_found = false} : !transform.any_op
   transform.annotate %1 "annotated" : !transform.any_op
 }
 
 // -----
 
 // CHECK-LABEL: func @patterns_apply_only_to_target_body()
 //       CHECK:   %0 = "test.foo"() {replace_with_new_op = "test.bar"} : () -> i32
 func.func @patterns_apply_only_to_target_body() {
   %0 = "test.foo"() {replace_with_new_op = "test.bar"} : () -> (i32)
   return
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%arg1: !transform.any_op):
 %0 = transform.structured.match ops{["test.foo"]} in %arg1 : (!transform.any_op) -> !transform.any_op
-  transform.apply_patterns ["transform.test"] to %0 {} : !transform.any_op
+  transform.apply_patterns to %0 {
+    transform.apply_patterns.transform.test_patterns
+  } : !transform.any_op
 }
 
 // -----
 
 // CHECK-LABEL: func @erase_tracked_op()
 //       CHECK:   "test.container"() ({
 //  CHECK-NEXT:   ^bb0:
 //  CHECK-NEXT:   }) : () -> ()
 func.func @erase_tracked_op() {
   "test.container"() ({
     // expected-remark @below {{matched op}}
     %0 = "test.erase_op"() {replace_with_new_op = "test.foo"} : () -> (i32)
   }) : () -> ()
   return
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%arg1: !transform.any_op):
   %0 = transform.structured.match ops{["test.container"]} in %arg1 : (!transform.any_op) -> !transform.any_op
   %1 = transform.structured.match ops{["test.erase_op"]} in %arg1 : (!transform.any_op) -> !transform.any_op
   transform.test_print_remark_at_operand %1, "matched op" : !transform.any_op
-  transform.apply_patterns ["transform.test"] to %0 {} : !transform.any_op
+  transform.apply_patterns to %0 {
+    transform.apply_patterns.transform.test_patterns
+  } : !transform.any_op
   transform.test_print_remark_at_operand %1, "op was deleted" : !transform.any_op
 }
 
 // -----
 
 // CHECK-LABEL: func @canonicalization(
 //       CHECK:   %[[c5:.*]] = arith.constant 5 : index
 //       CHECK:   return %[[c5]]
 func.func @canonicalization(%t: tensor<5xf32>) -> index {
   %c0 = arith.constant 0 : index
   // expected-remark @below {{op was replaced}}
   %dim = tensor.dim %t, %c0 : tensor<5xf32>
   return %dim : index
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%arg1: !transform.any_op):
   %0 = transform.structured.match ops{["tensor.dim"]} in %arg1 : (!transform.any_op) -> !transform.any_op
   %1 = transform.structured.match ops{["func.func"]} in %arg1 : (!transform.any_op) -> !transform.any_op
-  transform.apply_patterns [] to %1 {
+  transform.apply_patterns to %1 {
     transform.apply_patterns.canonicalization
   } : !transform.any_op
   transform.test_print_remark_at_operand %0, "op was replaced" : !transform.any_op
 }

mlir/test/Dialect/Vector/transform-vector.mlir

%b = transform.bufferization.one_shot_bufferize
       {bufferize_function_boundaries = true, allow_return_allocs = true}
       : (!transform.any_op) -> !transform.any_op
 
   %f = transform.structured.match ops{["func.func"]} in %b
     : (!transform.any_op) -> !transform.any_op
 
   // TODO: group these lower-level controls into various properly named vector
   // lowering TD macros.
-  transform.apply_patterns [] to %f {
+  transform.apply_patterns to %f {
     transform.apply_patterns.vector.lower_contraction lowering_strategy = "outerproduct"
   } : !transform.any_op
 
-  transform.apply_patterns [] to %f {
+  transform.apply_patterns to %f {
     transform.apply_patterns.vector.apply_transfer_permutation_patterns
   } : !transform.any_op
 
-  transform.apply_patterns [] to %f {
+  transform.apply_patterns to %f {
     transform.apply_patterns.vector.lower_multi_reduction lowering_strategy = "innerparallel"
   } : !transform.any_op
 
-  transform.apply_patterns [] to %f {
+  transform.apply_patterns to %f {
     transform.apply_patterns.vector.split_transfer_full_partial split_transfer_strategy = "linalg-copy"
   } : !transform.any_op
 
-  transform.apply_patterns [] to %f {
+  transform.apply_patterns to %f {
     transform.apply_patterns.vector.transfer_to_scf max_transfer_rank = 1 full_unroll = true
   } : !transform.any_op
 
-  transform.apply_patterns [] to %f {
+  transform.apply_patterns to %f {
     transform.apply_patterns.vector.lower_transfer max_transfer_rank = 1
   } : !transform.any_op
 
-  transform.apply_patterns [] to %f {
+  transform.apply_patterns to %f {
     transform.apply_patterns.vector.lower_shape_cast
   } : !transform.any_op
 
-  transform.apply_patterns [] to %f {
+  transform.apply_patterns to %f {
     transform.apply_patterns.vector.lower_transpose lowering_strategy = "shuffle_1d"
   } : !transform.any_op
 }

mlir/test/Dialect/Vector/vector-broadcast-lowering-transforms.mlir

func.func @broadcast_stretch_in_middle(%arg0: vector<4x1x2xf32>) -> vector<4x3x2xf32> {
   return %0 : vector<4x3x2xf32>
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
   %f = transform.structured.match ops{["func.func"]} in %module_op 
     : (!transform.any_op) -> !transform.any_op
 
-  transform.apply_patterns [] to %f {
+  transform.apply_patterns to %f {
     transform.apply_patterns.vector.lower_broadcast
   } : !transform.any_op
 }

mlir/test/Dialect/Vector/vector-contract-matvec-transforms.mlir

func.func @redpar_vecmattrans2x2(%arg0: memref<vector<2x2xf32>>, %arg1: memref<vector<2xf32>>,
   %b = memref.load %arg2[] : memref<vector<2xf32>>
   %0 = vector.contract #redpar_vecmattrans_trait %x, %A, %b : vector<2xf32>, vector<2x2xf32> into vector<2xf32>
   memref.store %0, %arg2[] : memref<vector<2xf32>>
   return
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
-  transform.apply_patterns [] to %module_op {
+  transform.apply_patterns to %module_op {
     transform.apply_patterns.vector.lower_contraction lowering_strategy = "outerproduct"
   } : !transform.any_op
 }

mlir/test/Dialect/Vector/vector-contract-to-dot-transforms.mlir

func.func @contract_one_sided_unit_reduction_dim(%arg0 : vector<1x2xi32>, %arg1 : vector<2x2xi32>, %arg2 : vector<2xi32>) -> vector<2xi32> {
   return %res : vector<2xi32>
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
   %f = transform.structured.match ops{["func.func"]} in %module_op 
     : (!transform.any_op) -> !transform.any_op
 
-  transform.apply_patterns [] to %f {
+  transform.apply_patterns to %f {
     transform.apply_patterns.vector.lower_contraction lowering_strategy = "dot"
   } : !transform.any_op
 }

mlir/test/Dialect/Vector/vector-contract-to-matrix-intrinsics-transforms.mlir

func.func @matmul(%arg0: vector<2x4xf32>,
   return %0 : vector<2x3xf32>
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
   %f = transform.structured.match ops{["func.func"]} in %module_op 
     : (!transform.any_op) -> !transform.any_op
 
-  transform.apply_patterns [] to %f {
+  transform.apply_patterns to %f {
     transform.apply_patterns.vector.lower_contraction lowering_strategy = "matmulintrinsics"
   } : !transform.any_op
 
-  transform.apply_patterns [] to %f {
+  transform.apply_patterns to %f {
     transform.apply_patterns.vector.lower_shape_cast
   } : !transform.any_op
 }

mlir/test/Dialect/Vector/vector-contract-to-outerproduct-transforms.mlir

 }
 
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
   %f = transform.structured.match ops{["func.func"]} in %module_op 
     : (!transform.any_op) -> !transform.any_op
 
-  transform.apply_patterns [] to %f {
+  transform.apply_patterns to %f {
     transform.apply_patterns.vector.lower_contraction lowering_strategy = "outerproduct"
   } : !transform.any_op
 }

mlir/test/Dialect/Vector/vector-contract-to-parallel-arith-transforms.mlir

func.func @parallel_contract_lowering_scalar(%arg0: vector<1x1xf32>, %arg1: vector<1x1xf32>, %arg2: f32) -> f32 {
   return %0 : f32
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
   %f = transform.structured.match ops{["func.func"]} in %module_op 
     : (!transform.any_op) -> !transform.any_op
 
-  transform.apply_patterns [] to %f {
+  transform.apply_patterns to %f {
     transform.apply_patterns.vector.lower_contraction lowering_strategy = "parallelarith"
   } : !transform.any_op
 }

mlir/test/Dialect/Vector/vector-mask-lowering-transforms.mlir

func.func @genbool_var_3d(%arg0: index, %arg1: index, %arg2: index) -> vector<2x1x7xi1> {
   return %0 : vector<2x1x7xi1>
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
   %f = transform.structured.match ops{["func.func"]} in %module_op 
     : (!transform.any_op) -> !transform.any_op
 
-  transform.apply_patterns [] to %f {
+  transform.apply_patterns to %f {
     transform.apply_patterns.vector.lower_masks
   } : !transform.any_op
 }
 
 // -----
 
 // CHECK-LABEL: func @transfer_read_3d(
 func.func @transfer_read_3d(
func.func @transfer_read_3d(
   return %1: vector<2x1x7xf32>
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
   %f = transform.structured.match ops{["func.func"]} in %module_op 
     : (!transform.any_op) -> !transform.any_op
 
-  transform.apply_patterns [] to %f {
+  transform.apply_patterns to %f {
     transform.apply_patterns.vector.lower_masked_transfers
   } : !transform.any_op
 }

mlir/test/Dialect/Vector/vector-multi-reduction-lowering.mlir

 }
 
 // CHECK-LABEL: func @vector_multi_reduction_parallel_middle
 //  CHECK-SAME:   %[[INPUT:.+]]: vector<3x4x5xf32>, %[[ACC:.+]]: vector<4xf32>
 //       CHECK: vector.transpose %[[INPUT]], [1, 0, 2] : vector<3x4x5xf32> to vector<4x3x5xf32>
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
-  transform.apply_patterns [] to %module_op {
+  transform.apply_patterns to %module_op {
     transform.apply_patterns.vector.lower_multi_reduction lowering_strategy = "innerreduction"
   } : !transform.any_op
 }

mlir/test/Dialect/Vector/vector-multi-reduction-outer-lowering.mlir

func.func @vector_multi_reduction_to_scalar(%arg0: vector<2x3xf32>, %acc: f32) -> f32 {
   %0 = vector.multi_reduction <add>, %arg0, %acc [0, 1] : vector<2x3xf32> to f32
   return %0 : f32
 }
 // CHECK-LABEL: func @vector_multi_reduction_to_scalar
 //       CHECK:   return %{{.+}}
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
-  transform.apply_patterns [] to %module_op {
+  transform.apply_patterns to %module_op {
     transform.apply_patterns.vector.lower_multi_reduction lowering_strategy = "innerparallel"
   } : !transform.any_op
 }

mlir/test/Dialect/Vector/vector-outerproduct-lowering-transforms.mlir

func.func @axpy_int_add(%arg0: vector<16xi32>, %arg1: i32, %arg2: vector<16xi32>) -> vector<16xi32> {
   return %0: vector<16xi32>
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
   %f = transform.structured.match ops{["func.func"]} in %module_op 
     : (!transform.any_op) -> !transform.any_op
 
-  transform.apply_patterns [] to %f {
+  transform.apply_patterns to %f {
     transform.apply_patterns.vector.lower_outerproduct
   } : !transform.any_op
 
-  transform.apply_patterns [] to %f {
+  transform.apply_patterns to %f {
     transform.apply_patterns.vector.lower_broadcast
   } : !transform.any_op
 }

mlir/test/Dialect/Vector/vector-shape-cast-lowering-transforms.mlir

func.func @shape_cast_1d0d(%arg0 : vector<1xf32>) -> vector<f32> {
   return %s : vector<f32>
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
   %f = transform.structured.match ops{["func.func"]} in %module_op
     : (!transform.any_op) -> !transform.any_op
 
-  transform.apply_patterns [] to %f {
+  transform.apply_patterns to %f {
     transform.apply_patterns.vector.lower_shape_cast
   } : !transform.any_op
 }

mlir/test/Dialect/Vector/vector-transfer-drop-unit-dims-patterns.mlir

 // CHECK-LABEL: func @transfer_read_rank_reducing
 //  CHECK-SAME:     %[[ARG:.+]]: memref<1x1x3x2xi8
 //       CHECK:   %[[SUBVIEW:.+]] = memref.subview %[[ARG]][0, 0, 0, 0] [1, 1, 3, 2] [1, 1, 1, 1]
 //  CHECK-SAME:     memref<1x1x3x2xi8, {{.*}}> to memref<3x2xi8, {{.*}}>
 //       CHECK:   vector.transfer_read %[[SUBVIEW]]
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
-  transform.apply_patterns [] to %module_op {
+  transform.apply_patterns to %module_op {
     transform.apply_patterns.vector.apply_rank_reducing_subview_patterns
   } : !transform.any_op
 }
 
 // -----
 
 func.func @transfer_write_rank_reducing(%arg : memref<1x1x3x2xi8, strided<[6, 6, 2, 1], offset: ?>>, %vec : vector<3x2xi8>) {
     %c0 = arith.constant 0 : index
     vector.transfer_write %vec, %arg [%c0, %c0, %c0, %c0] :
       vector<3x2xi8>, memref<1x1x3x2xi8, strided<[6, 6, 2, 1], offset: ?>>
     return
 }
 
 // CHECK-LABEL: func @transfer_write_rank_reducing
 //  CHECK-SAME:     %[[ARG:.+]]: memref<1x1x3x2xi8
 //       CHECK:   %[[SUBVIEW:.+]] = memref.subview %[[ARG]][0, 0, 0, 0] [1, 1, 3, 2] [1, 1, 1, 1]
 //  CHECK-SAME:     memref<1x1x3x2xi8, {{.*}}> to memref<3x2xi8, {{.*}}>
 //       CHECK:   vector.transfer_write %{{.*}}, %[[SUBVIEW]]
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
-  transform.apply_patterns [] to %module_op {
+  transform.apply_patterns to %module_op {
     transform.apply_patterns.vector.apply_rank_reducing_subview_patterns
   } : !transform.any_op
 }
 
 // -----
 
 func.func @transfer_read_and_vector_rank_reducing(
       %arg : memref<1x1x3x2x1xf32>) -> vector<3x2x1xf32> {
     %c0 = arith.constant 0 : index
     %cst = arith.constant 0.0 : f32
     %v = vector.transfer_read %arg[%c0, %c0, %c0, %c0, %c0], %cst :
       memref<1x1x3x2x1xf32>, vector<3x2x1xf32>
     return %v : vector<3x2x1xf32>
 }
 
 // CHECK-LABEL: func @transfer_read_and_vector_rank_reducing
 //  CHECK-SAME:     %[[ARG:.+]]: memref<1x1x3x2x1xf32>
 //       CHECK:   %[[SUBVIEW:.+]] = memref.subview %[[ARG]][0, 0, 0, 0, 0] [1, 1, 3, 2, 1] [1, 1, 1, 1, 1]
 //  CHECK-SAME:     memref<1x1x3x2x1xf32> to memref<3x2xf32>
 //       CHECK:   vector.transfer_read %[[SUBVIEW]]{{.*}} {in_bounds = [true, true]} : memref<3x2xf32>, vector<3x2xf32>
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
-  transform.apply_patterns [] to %module_op {
+  transform.apply_patterns to %module_op {
     transform.apply_patterns.vector.apply_rank_reducing_subview_patterns
   } : !transform.any_op
 }
 
 // -----
 
 func.func @transfer_write_and_vector_rank_reducing(
       %arg : memref<1x1x3x2x1xf32>,
       %vec : vector<3x2x1xf32>) {
     %c0 = arith.constant 0 : index
     vector.transfer_write %vec, %arg [%c0, %c0, %c0, %c0, %c0] :
       vector<3x2x1xf32>, memref<1x1x3x2x1xf32>
     return
 }
 
 // CHECK-LABEL: func @transfer_write_and_vector_rank_reducing
 //  CHECK-SAME:     %[[ARG:.+]]: memref<1x1x3x2x1xf32>
 //       CHECK:   %[[SUBVIEW:.+]] = memref.subview %[[ARG]][0, 0, 0, 0, 0] [1, 1, 3, 2, 1] [1, 1, 1, 1, 1]
 //  CHECK-SAME:     memref<1x1x3x2x1xf32> to memref<3x2xf32>
 //       CHECK:   vector.transfer_write %{{.*}}, %[[SUBVIEW]]{{.*}} {in_bounds = [true, true]} : vector<3x2xf32>, memref<3x2xf32>
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
-  transform.apply_patterns [] to %module_op {
+  transform.apply_patterns to %module_op {
     transform.apply_patterns.vector.apply_rank_reducing_subview_patterns
   } : !transform.any_op
 }
 
 // -----
 
 func.func @transfer_read_and_vector_rank_reducing_to_0d(
       %arg : memref<1x1x1x1x1xf32>) -> vector<1x1x1xf32> {
     %c0 = arith.constant 0 : index
     %cst = arith.constant 0.0 : f32
     %v = vector.transfer_read %arg[%c0, %c0, %c0, %c0, %c0], %cst :
       memref<1x1x1x1x1xf32>, vector<1x1x1xf32>
     return %v : vector<1x1x1xf32>
 }
 
 // CHECK-LABEL: func @transfer_read_and_vector_rank_reducing_to_0d
 //  CHECK-SAME:     %[[MEMREF:.+]]: memref<1x1x1x1x1xf32>
 //       CHECK:   %[[SUBVIEW:.+]] = memref.subview %[[MEMREF]][0, 0, 0, 0, 0] [1, 1, 1, 1, 1] [1, 1, 1, 1, 1] : memref<1x1x1x1x1xf32> to memref<f32>
 //       CHECK:   %[[READ:.+]] = vector.transfer_read %[[SUBVIEW]]{{.*}} : memref<f32>, vector<f32>
 //       CHECK:   vector.shape_cast %[[READ]] : vector<f32> to vector<1x1x1xf32>
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
-  transform.apply_patterns [] to %module_op {
+  transform.apply_patterns to %module_op {
     transform.apply_patterns.vector.apply_rank_reducing_subview_patterns
   } : !transform.any_op
 }
 
 // -----
 
 func.func @transfer_write_and_vector_rank_reducing_to_0d(
       %arg : memref<1x1x1x1x1xf32>,
       %vec : vector<1x1x1xf32>) {
     %c0 = arith.constant 0 : index
     vector.transfer_write %vec, %arg [%c0, %c0, %c0, %c0, %c0] :
       vector<1x1x1xf32>, memref<1x1x1x1x1xf32>
     return
 }
 
 // CHECK-LABEL: func @transfer_write_and_vector_rank_reducing_to_0d
 //  CHECK-SAME:     %[[MEMREF:.+]]: memref<1x1x1x1x1xf32>, %[[VECTOR:.+]]: vector<1x1x1xf32>
 //       CHECK:   %[[SUBVIEW:.+]] = memref.subview %[[MEMREF]][0, 0, 0, 0, 0] [1, 1, 1, 1, 1] [1, 1, 1, 1, 1] : memref<1x1x1x1x1xf32> to memref<f32>
 //       CHECK:   %[[SHCAST:.+]] = vector.shape_cast %[[VECTOR]] : vector<1x1x1xf32> to vector<f32>
 //       CHECK:   vector.transfer_write %[[SHCAST]], %[[SUBVIEW]]{{.*}} : vector<f32>, memref<f32>
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
-  transform.apply_patterns [] to %module_op {
+  transform.apply_patterns to %module_op {
     transform.apply_patterns.vector.apply_rank_reducing_subview_patterns
   } : !transform.any_op
 }

mlir/test/Dialect/Vector/vector-transfer-full-partial-split-copy-transform.mlir

func.func @split_vector_transfer_read_strided_2d(
   %1 = vector.transfer_read %A[%i, %j], %f0 :
     memref<7x8xf32, strided<[?, 1], offset: ?>>, vector<4x8xf32>
 
   return %1 : vector<4x8xf32>
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
-  transform.apply_patterns [] to %module_op {
+  transform.apply_patterns to %module_op {
     transform.apply_patterns.vector.split_transfer_full_partial split_transfer_strategy = "linalg-copy"
   } : !transform.any_op
 }
 
 // -----
 
 func.func @split_vector_transfer_write_2d(%V: vector<4x8xf32>, %A: memref<?x8xf32>, %i: index, %j: index) {
   vector.transfer_write %V, %A[%i, %j] :

 // CHECK:             memref.copy %[[VAL_22]], %[[DEST_VIEW]]
 // CHECK-SAME:            : memref<?x?xf32, strided<[8, 1], offset: ?>> to memref<?x?xf32, strided{{.*}}>
 // CHECK:           }
 // CHECK:           return
 // CHECK:         }
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
-  transform.apply_patterns [] to %module_op {
+  transform.apply_patterns to %module_op {
     transform.apply_patterns.vector.split_transfer_full_partial split_transfer_strategy = "linalg-copy"
   } : !transform.any_op
 }
 
 // -----
 
 func.func @split_vector_transfer_write_strided_2d(
     %V: vector<4x8xf32>, %A: memref<7x8xf32, strided<[?, 1], offset: ?>>,

 // CHECK:             memref.copy %[[VAL_22]], %[[DEST_VIEW]]
 // CHECK-SAME:            : memref<?x?xf32, strided<[8, 1], offset: ?>> to memref<?x?xf32, strided<[?, 1], offset: ?>>
 // CHECK:           }
 // CHECK:           return
 // CHECK:         }
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
-  transform.apply_patterns [] to %module_op {
+  transform.apply_patterns to %module_op {
     transform.apply_patterns.vector.split_transfer_full_partial split_transfer_strategy = "linalg-copy"
   } : !transform.any_op
 }

mlir/test/Dialect/Vector/vector-transfer-full-partial-split.mlir

%1 = vector.transfer_read %A[%i, %j], %f0 :
     memref<7x8xf32, strided<[?, 1], offset: ?>>, vector<4x8xf32>
 
   // CHECK: return %[[res]] : vector<4x8xf32>
   return %1 : vector<4x8xf32>
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
-  transform.apply_patterns [] to %module_op {
+  transform.apply_patterns to %module_op {
     transform.apply_patterns.vector.split_transfer_full_partial split_transfer_strategy = "vector-transfer"
   } : !transform.any_op
 }
 
 // -----
 
 func.func @split_vector_transfer_write_2d(%V: vector<4x8xf32>, %A: memref<?x8xf32>, %i: index, %j: index) {
   vector.transfer_write %V, %A[%i, %j] :

 // CHECK-SAME:            : vector<4x8xf32>, memref<?x8xf32>
 // CHECK:           }
 // CHECK:           return
 // CHECK:         }
 
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
-  transform.apply_patterns [] to %module_op {
+  transform.apply_patterns to %module_op {
     transform.apply_patterns.vector.split_transfer_full_partial split_transfer_strategy = "vector-transfer"
   } : !transform.any_op
 }
 
 // -----
 
 func.func @split_vector_transfer_write_strided_2d(
     %V: vector<4x8xf32>, %A: memref<7x8xf32, strided<[?, 1], offset: ?>>,

 // CHECK:             vector.transfer_write %[[VAL_20]], %[[DEST]][%[[I]], %[[J]]]
 // CHECK-SAME:            : vector<4x8xf32>, memref<7x8xf32, strided<[?, 1], offset: ?>>
 // CHECK:           }
 // CHECK:           return
 // CHECK:         }
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
-  transform.apply_patterns [] to %module_op {
+  transform.apply_patterns to %module_op {
     transform.apply_patterns.vector.split_transfer_full_partial split_transfer_strategy = "vector-transfer"
   } : !transform.any_op
 }
 
 // -----
 
 func.func private @fake_side_effecting_fun(%0: vector<2x2xf32>) -> ()
 
scf.for %i = %lb to %ub step %step {
     %0 = vector.transfer_read %A[%c0, %c0], %f0 : memref<?x?xf32>, vector<2x2xf32>
     func.call @fake_side_effecting_fun(%0) : (vector<2x2xf32>) -> ()
   }
   return
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
-  transform.apply_patterns [] to %module_op {
+  transform.apply_patterns to %module_op {
     transform.apply_patterns.vector.split_transfer_full_partial split_transfer_strategy = "vector-transfer"
   } : !transform.any_op
 }

mlir/test/Dialect/Vector/vector-transfer-to-vector-load-store.mlir

%res = vector.transfer_read %mem[%i, %i, %i, %i, %i], %cf0
     {in_bounds = [true, true, true, true], permutation_map = #broadcast_2d_in_4d} 
       : memref<10x20x30x8x8xf32>, vector<3x2x4x5xf32>
   return %res : vector<3x2x4x5xf32>
 }
 
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
-  transform.apply_patterns [] to %module_op {
+  transform.apply_patterns to %module_op {
     transform.apply_patterns.vector.lower_transfer max_transfer_rank = 99
     transform.apply_patterns.vector.apply_transfer_permutation_patterns
   } : !transform.any_op
 }
 
 // -----
 
 #map0 = affine_map<(d0, d1, d2, d3) -> (d1, d0, 0, 0)>
func.func @transfer_write_broadcast_unit_dim(
   // CHECK: %[[NEW_VEC3:.*]] = vector.transpose %[[NEW_VEC2]], [1, 2, 0] : vector<1x8x16xf32> to vector<8x16x1xf32>
   // CHECK: vector.transfer_write %[[NEW_VEC3]], %[[ARG0]][%[[C0]], %[[C0]], %[[C0]], %[[C0]]] : vector<8x16x1xf32>, memref<?x?x?x?xf32>
 
   return %0 : tensor<?x?x?x?xf32>
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
-  transform.apply_patterns [] to %module_op {
+  transform.apply_patterns to %module_op {
     transform.apply_patterns.vector.lower_transfer max_transfer_rank = 99
     transform.apply_patterns.vector.apply_transfer_permutation_patterns
   } : !transform.any_op
 }

mlir/test/Dialect/Vector/vector-transpose-lowering.mlir

func.func @transpose1023_1x1x8x8xf32(%arg0: vector<1x1x8x8xf32>) -> vector<1x1x8x8xf32> {
   //      CHECK: vector.extract {{.*}}[0, 0] : vector<1x1x8x8xf32>
   // CHECK-NEXT: vector.insert {{.*}} [0, 0] : vector<8x8xf32> into vector<1x1x8x8xf32>
   %0 = vector.transpose %arg0, [1, 0, 2, 3] : vector<1x1x8x8xf32> to vector<1x1x8x8xf32>
   return %0 : vector<1x1x8x8xf32>
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
-  transform.apply_patterns [] to %module_op {
+  transform.apply_patterns to %module_op {
     transform.apply_patterns.vector.lower_transpose lowering_strategy = "eltwise"
   } : !transform.any_op
 }
 
 // -----
 
 // CHECK-LABEL: func @transpose
 func.func @transpose(%arg0: vector<2x4xf32>) -> vector<4x2xf32> {
   //      CHECK: vector.shape_cast %{{.*}} : vector<2x4xf32> to vector<8xf32>
   //            0 4
   // 0 1 2 3    1 5
   // 4 5 6 7 -> 2 6
   //            3 7
   // CHECK-NEXT: vector.shuffle %{{.*}} [0, 4, 1, 5, 2, 6, 3, 7] : vector<8xf32>, vector<8xf32>
   // CHECK-NEXT: vector.shape_cast %{{.*}} : vector<8xf32> to vector<4x2xf32>
   %0 = vector.transpose %arg0, [1, 0] : vector<2x4xf32> to vector<4x2xf32>
   return %0 : vector<4x2xf32>
 }
 
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
-  transform.apply_patterns [] to %module_op {
+  transform.apply_patterns to %module_op {
     transform.apply_patterns.vector.lower_transpose lowering_strategy = "shuffle_1d"
   } : !transform.any_op
 }
 
 // -----
 
 // CHECK-LABEL: func @transpose(
 func.func @transpose(%arg0: vector<2x4xf32>) -> vector<4x2xf32> {
   // CHECK:       vector.shape_cast {{.*}} : vector<2x4xf32> to vector<8xf32>
   // CHECK:       vector.flat_transpose %{{.*}} {columns = 2 : i32, rows = 4 : i32} : vector<8xf32> -> vector<8xf32>
   // CHECK:       vector.shape_cast {{.*}} : vector<8xf32> to vector<4x2xf32>
   %0 = vector.transpose %arg0, [1, 0] : vector<2x4xf32> to vector<4x2xf32>
   return %0 : vector<4x2xf32>
 }
 
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
-  transform.apply_patterns [] to %module_op {
+  transform.apply_patterns to %module_op {
     transform.apply_patterns.vector.lower_transpose lowering_strategy = "flat_transpose"
   } : !transform.any_op
 }
 
 // -----
 
 // CHECK-LABEL: func @transpose4x8
 func.func @transpose4x8xf32(%arg0: vector<4x8xf32>) -> vector<8x4xf32> {
func.func @transpose210_1x8x8xf32(%arg0: vector<1x8x8xf32>) -> vector<8x8x1xf32> {
   // CHECK-NEXT: vector.insert {{.*}}[7]
   // CHECK-NEXT: vector.shape_cast %{{.*}} : vector<8x8xf32> to vector<8x8x1xf32>
   %0 = vector.transpose %arg0, [2, 1, 0] : vector<1x8x8xf32> to vector<8x8x1xf32>
   return %0 : vector<8x8x1xf32>
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
-  transform.apply_patterns [] to %module_op {
+  transform.apply_patterns to %module_op {
     transform.apply_patterns.vector.lower_transpose avx2_lowering_strategy = true
   } : !transform.any_op
 }
 
 // -----
 
 func.func @transpose_shuffle16x16xf32(%arg0: vector<16x16xf32>) -> vector<16x16xf32> {
   // CHECK: vector.shuffle {{.*}} [0, 16, 1, 17, 4, 20, 5, 21, 8, 24, 9, 25, 12, 28, 13, 29] : vector<16xf32>, vector<16xf32>
func.func @transpose_shuffle16x16xf32(%arg0: vector<16x16xf32>) -> vector<16x16xf32> {
   // CHECK: vector.shuffle {{.*}} [4, 5, 6, 7, 12, 13, 14, 15, 20, 21, 22, 23, 28, 29, 30, 31] : vector<16xf32>, vector<16xf32>
   // CHECK: vector.shuffle {{.*}} [4, 5, 6, 7, 12, 13, 14, 15, 20, 21, 22, 23, 28, 29, 30, 31] : vector<16xf32>, vector<16xf32>
   %0 = vector.transpose %arg0, [1, 0] : vector<16x16xf32> to vector<16x16xf32>
   return %0 : vector<16x16xf32>
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
-  transform.apply_patterns [] to %module_op {
+  transform.apply_patterns to %module_op {
     transform.apply_patterns.vector.lower_transpose lowering_strategy = "shuffle_16x16"
   } : !transform.any_op
 }
 
 // -----
 
 // CHECK-LABEL: func @transpose021_shuffle16x16xf32
 func.func @transpose021_shuffle16x16xf32(%arg0: vector<1x16x16xf32>) -> vector<1x16x16xf32> {
func.func @transpose021_shuffle16x16xf32(%arg0: vector<1x16x16xf32>) -> vector<1x16x16xf32> {
   // CHECK: vector.shuffle {{.*}} [4, 5, 6, 7, 12, 13, 14, 15, 20, 21, 22, 23, 28, 29, 30, 31] : vector<16xf32>, vector<16xf32>
   // CHECK: vector.shuffle {{.*}} [4, 5, 6, 7, 12, 13, 14, 15, 20, 21, 22, 23, 28, 29, 30, 31] : vector<16xf32>, vector<16xf32>
   %0 = vector.transpose %arg0, [0, 2, 1] : vector<1x16x16xf32> to vector<1x16x16xf32>
   return %0 : vector<1x16x16xf32>
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%module_op: !transform.any_op):
-  transform.apply_patterns [] to %module_op {
+  transform.apply_patterns to %module_op {
     transform.apply_patterns.vector.lower_transpose lowering_strategy = "shuffle_16x16"
   } : !transform.any_op
 }

mlir/test/lib/Dialect/Transform/TestTransformDialectExtension.cpp

 } // namespace
 
 void mlir::test::ApplyTestPatternsOp::populatePatterns(
     RewritePatternSet &patterns) {
   patterns.insert<ReplaceWithNewOp, EraseOp>(patterns.getContext());
 }
 
 namespace {
-void populateTestPatterns(RewritePatternSet &patterns) {
-  patterns.insert<ReplaceWithNewOp, EraseOp>(patterns.getContext());
-}
-
 /// Test extension of the Transform dialect. Registers additional ops and
 /// declares PDL as dependent dialect since the additional ops are using PDL
 /// types for operands and results.
 class TestTransformDialectExtension
     : public transform::TransformDialectExtension<
           TestTransformDialectExtension> {
 public:
   using Base::Base;
#include "TestTransformDialectExtensionTypes.cpp.inc"
     };
 
     addDialectDataInitializer<transform::PDLMatchHooks>(
         [&](transform::PDLMatchHooks &hooks) {
           llvm::StringMap<PDLConstraintFunction> constraints;
           constraints.try_emplace("verbose_constraint", verboseConstraint);
           hooks.mergeInPDLMatchHooks(std::move(constraints));
         });
-
-    addDialectDataInitializer<transform::PatternRegistry>(
-        [&](transform::PatternRegistry &registry) {
-          registry.registerPatterns("transform.test", populateTestPatterns);
-        });
   }
 };
 } // namespace
 
 // These are automatically generated by ODS but are not used as the Transform
 // dialect uses a different dispatch mechanism to support dialect extensions.
 LLVM_ATTRIBUTE_UNUSED static OptionalParseResult
 generatedTypeParser(AsmParser &parser, StringRef *mnemonic, Type &value);