[Mlir-commits] [mlir] cc7f524 - [mlir][transform] Use separate ops instead of PatternRegistry
Matthias Springer
llvmlistbot at llvm.org
Tue Jun 6 02:54:39 PDT 2023
Author: Matthias Springer
Date: 2023-06-06T11:53:03+02:00
New Revision: cc7f52432bca6938d748c6730943586f879f841f
URL: https://github.com/llvm/llvm-project/commit/cc7f52432bca6938d748c6730943586f879f841f
DIFF: https://github.com/llvm/llvm-project/commit/cc7f52432bca6938d748c6730943586f879f841f.diff
LOG: [mlir][transform] Use separate ops instead of PatternRegistry
* 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.
Differential Revision: https://reviews.llvm.org/D152249
Added:
Modified:
mlir/include/mlir/Dialect/Linalg/TransformOps/LinalgTransformOps.td
mlir/include/mlir/Dialect/MemRef/TransformOps/MemRefTransformOps.td
mlir/include/mlir/Dialect/SCF/TransformOps/SCFTransformOps.td
mlir/include/mlir/Dialect/Tensor/TransformOps/TensorTransformOps.h
mlir/include/mlir/Dialect/Tensor/TransformOps/TensorTransformOps.td
mlir/include/mlir/Dialect/Tensor/Transforms/Transforms.h
mlir/include/mlir/Dialect/Transform/IR/TransformOps.h
mlir/include/mlir/Dialect/Transform/IR/TransformOps.td
mlir/include/mlir/Dialect/Vector/TransformOps/VectorTransformOps.td
mlir/lib/Dialect/Linalg/TransformOps/DialectExtension.cpp
mlir/lib/Dialect/Linalg/TransformOps/LinalgTransformOps.cpp
mlir/lib/Dialect/MemRef/TransformOps/MemRefTransformOps.cpp
mlir/lib/Dialect/SCF/TransformOps/SCFTransformOps.cpp
mlir/lib/Dialect/Tensor/TransformOps/TensorTransformOps.cpp
mlir/lib/Dialect/Transform/IR/TransformOps.cpp
mlir/lib/Dialect/Vector/TransformOps/VectorTransformOps.cpp
mlir/test/Dialect/LLVM/transform-e2e.mlir
mlir/test/Dialect/Linalg/transform-op-matmul-to-outerproduct.mlir
mlir/test/Dialect/MemRef/extract-address-computations.mlir
mlir/test/Dialect/Transform/ops-invalid.mlir
mlir/test/Dialect/Transform/test-pattern-application.mlir
mlir/test/Dialect/Vector/transform-vector.mlir
mlir/test/Dialect/Vector/vector-broadcast-lowering-transforms.mlir
mlir/test/Dialect/Vector/vector-contract-matvec-transforms.mlir
mlir/test/Dialect/Vector/vector-contract-to-dot-transforms.mlir
mlir/test/Dialect/Vector/vector-contract-to-matrix-intrinsics-transforms.mlir
mlir/test/Dialect/Vector/vector-contract-to-outerproduct-transforms.mlir
mlir/test/Dialect/Vector/vector-contract-to-parallel-arith-transforms.mlir
mlir/test/Dialect/Vector/vector-mask-lowering-transforms.mlir
mlir/test/Dialect/Vector/vector-multi-reduction-lowering.mlir
mlir/test/Dialect/Vector/vector-multi-reduction-outer-lowering.mlir
mlir/test/Dialect/Vector/vector-outerproduct-lowering-transforms.mlir
mlir/test/Dialect/Vector/vector-shape-cast-lowering-transforms.mlir
mlir/test/Dialect/Vector/vector-transfer-drop-unit-dims-patterns.mlir
mlir/test/Dialect/Vector/vector-transfer-full-partial-split-copy-transform.mlir
mlir/test/Dialect/Vector/vector-transfer-full-partial-split.mlir
mlir/test/Dialect/Vector/vector-transfer-to-vector-load-store.mlir
mlir/test/Dialect/Vector/vector-transpose-lowering.mlir
mlir/test/Integration/Dialect/Vector/CPU/test-shuffle16x16.mlir
mlir/test/lib/Dialect/Transform/TestTransformDialectExtension.cpp
Removed:
################################################################################
diff --git a/mlir/include/mlir/Dialect/Linalg/TransformOps/LinalgTransformOps.td b/mlir/include/mlir/Dialect/Linalg/TransformOps/LinalgTransformOps.td
index 65a2510a11bca..a7356c0e57e20 100644
--- a/mlir/include/mlir/Dialect/Linalg/TransformOps/LinalgTransformOps.td
+++ b/mlir/include/mlir/Dialect/Linalg/TransformOps/LinalgTransformOps.td
@@ -26,6 +26,53 @@ def TransformParamTypeOrAnyHandle : Type<
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
//===----------------------------------------------------------------------===//
diff --git a/mlir/include/mlir/Dialect/MemRef/TransformOps/MemRefTransformOps.td b/mlir/include/mlir/Dialect/MemRef/TransformOps/MemRefTransformOps.td
index 04dfe1f430241..35a1d840f2d69 100644
--- a/mlir/include/mlir/Dialect/MemRef/TransformOps/MemRefTransformOps.td
+++ b/mlir/include/mlir/Dialect/MemRef/TransformOps/MemRefTransformOps.td
@@ -15,6 +15,81 @@ 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",
@@ -48,51 +123,6 @@ def MemRefMultiBufferOp : Op<Transform_Dialect, "memref.multibuffer",
"$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,
diff --git a/mlir/include/mlir/Dialect/SCF/TransformOps/SCFTransformOps.td b/mlir/include/mlir/Dialect/SCF/TransformOps/SCFTransformOps.td
index fef9e4bd17214..d128c0e4f2796 100644
--- a/mlir/include/mlir/Dialect/SCF/TransformOps/SCFTransformOps.td
+++ b/mlir/include/mlir/Dialect/SCF/TransformOps/SCFTransformOps.td
@@ -15,6 +15,18 @@ 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",
diff --git a/mlir/include/mlir/Dialect/Tensor/TransformOps/TensorTransformOps.h b/mlir/include/mlir/Dialect/Tensor/TransformOps/TensorTransformOps.h
index d3815a3401a21..0e5fda6041faa 100644
--- a/mlir/include/mlir/Dialect/Tensor/TransformOps/TensorTransformOps.h
+++ b/mlir/include/mlir/Dialect/Tensor/TransformOps/TensorTransformOps.h
@@ -9,7 +9,7 @@
#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"
diff --git a/mlir/include/mlir/Dialect/Tensor/TransformOps/TensorTransformOps.td b/mlir/include/mlir/Dialect/Tensor/TransformOps/TensorTransformOps.td
index a9580bda39420..22367c98fdc72 100644
--- a/mlir/include/mlir/Dialect/Tensor/TransformOps/TensorTransformOps.td
+++ b/mlir/include/mlir/Dialect/Tensor/TransformOps/TensorTransformOps.td
@@ -15,6 +15,75 @@ 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
diff --git a/mlir/include/mlir/Dialect/Tensor/Transforms/Transforms.h b/mlir/include/mlir/Dialect/Tensor/Transforms/Transforms.h
index a87ee1b3e8053..cadf3ad55f6a1 100644
--- a/mlir/include/mlir/Dialect/Tensor/Transforms/Transforms.h
+++ b/mlir/include/mlir/Dialect/Tensor/Transforms/Transforms.h
@@ -39,8 +39,8 @@ FailureOr<TilingResult> replaceExtractSliceWithTiledProducer(
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);
@@ -55,7 +55,7 @@ void populateDropRedundantInsertSliceRankExpansionPatterns(
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`
diff --git a/mlir/include/mlir/Dialect/Transform/IR/TransformOps.h b/mlir/include/mlir/Dialect/Transform/IR/TransformOps.h
index e738baf15c8f9..5625d193b1508 100644
--- a/mlir/include/mlir/Dialect/Transform/IR/TransformOps.h
+++ b/mlir/include/mlir/Dialect/Transform/IR/TransformOps.h
@@ -26,7 +26,6 @@
namespace mlir {
namespace transform {
-class ApplyPatternsOp;
enum class FailurePropagationMode : uint32_t;
class FailurePropagationModeAttr;
@@ -152,51 +151,9 @@ class ErrorCheckingTrackingListener : public TrackingListener {
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"
diff --git a/mlir/include/mlir/Dialect/Transform/IR/TransformOps.td b/mlir/include/mlir/Dialect/Transform/IR/TransformOps.td
index 5b7e6ca8ca14e..6c75e4b622682 100644
--- a/mlir/include/mlir/Dialect/Transform/IR/TransformOps.td
+++ b/mlir/include/mlir/Dialect/Transform/IR/TransformOps.td
@@ -138,12 +138,9 @@ def ApplyPatternsOp : TransformDialectOp<"apply_patterns",
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.
@@ -161,12 +158,12 @@ def ApplyPatternsOp : TransformDialectOp<"apply_patterns",
}];
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 = [{
diff --git a/mlir/include/mlir/Dialect/Vector/TransformOps/VectorTransformOps.td b/mlir/include/mlir/Dialect/Vector/TransformOps/VectorTransformOps.td
index 87fcbb0127e7f..c1891ff4463c9 100644
--- a/mlir/include/mlir/Dialect/Vector/TransformOps/VectorTransformOps.td
+++ b/mlir/include/mlir/Dialect/Vector/TransformOps/VectorTransformOps.td
@@ -52,7 +52,7 @@ def ApplyTransferPermutationPatternsOp : Op<Transform_Dialect,
let assemblyFormat = "attr-dict";
}
-def LowerBroadcastOp : Op<Transform_Dialect,
+def ApplyLowerBroadcastPatternsOp : Op<Transform_Dialect,
"apply_patterns.vector.lower_broadcast",
[DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
let description = [{
@@ -67,7 +67,7 @@ def LowerBroadcastOp : Op<Transform_Dialect,
}
// 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 = [{
@@ -86,7 +86,7 @@ def LowerContractionOp : Op<Transform_Dialect,
}];
}
-def LowerMasksOp : Op<Transform_Dialect,
+def ApplyLowerMasksPatternsOp : Op<Transform_Dialect,
"apply_patterns.vector.lower_masks",
[DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
let description = [{
@@ -100,7 +100,7 @@ def LowerMasksOp : Op<Transform_Dialect,
let assemblyFormat = "attr-dict";
}
-def LowerMaskedTransfersOp : Op<Transform_Dialect,
+def ApplyLowerMaskedTransfersPatternsOp : Op<Transform_Dialect,
"apply_patterns.vector.lower_masked_transfers",
[DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
let description = [{
@@ -114,7 +114,7 @@ def LowerMaskedTransfersOp : Op<Transform_Dialect,
let assemblyFormat = "attr-dict";
}
-def MaterializeMasksOp : Op<Transform_Dialect,
+def ApplyMaterializeMasksPatternsOp : Op<Transform_Dialect,
"apply_patterns.vector.materialize_masks",
[DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
let description = [{
@@ -129,7 +129,7 @@ def MaterializeMasksOp : Op<Transform_Dialect,
}
// 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 = [{
@@ -149,7 +149,7 @@ def LowerMultiReductionOp : Op<Transform_Dialect,
}];
}
-def LowerOuterProductOp : Op<Transform_Dialect,
+def ApplyLowerOuterProductPatternsOp : Op<Transform_Dialect,
"apply_patterns.vector.lower_outerproduct",
[DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
let description = [{
@@ -163,7 +163,29 @@ def LowerOuterProductOp : Op<Transform_Dialect,
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 = [{
@@ -177,7 +199,7 @@ def LowerShapeCastOp : Op<Transform_Dialect,
let assemblyFormat = "attr-dict";
}
-def LowerTransferOp : Op<Transform_Dialect,
+def ApplyLowerTransferPatternsOp : Op<Transform_Dialect,
"apply_patterns.vector.lower_transfer",
[DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
let description = [{
@@ -196,7 +218,7 @@ def LowerTransferOp : Op<Transform_Dialect,
}
// 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 = [{
@@ -223,7 +245,7 @@ def LowerTransposeOp : Op<Transform_Dialect,
}
// 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 = [{
@@ -244,7 +266,7 @@ def SplitTransferFullPartialOp : Op<Transform_Dialect,
}];
}
-def TransferToScfOp : Op<Transform_Dialect,
+def ApplyTransferToScfPatternsOp : Op<Transform_Dialect,
"apply_patterns.vector.transfer_to_scf",
[DeclareOpInterfaceMethods<PatternDescriptorOpInterface>]> {
let description = [{
diff --git a/mlir/lib/Dialect/Linalg/TransformOps/DialectExtension.cpp b/mlir/lib/Dialect/Linalg/TransformOps/DialectExtension.cpp
index ab28ad32e350d..2bf2a569b3333 100644
--- a/mlir/lib/Dialect/Linalg/TransformOps/DialectExtension.cpp
+++ b/mlir/lib/Dialect/Linalg/TransformOps/DialectExtension.cpp
@@ -49,22 +49,6 @@ class LinalgTransformDialectExtension
#define GET_OP_LIST
#include "mlir/Dialect/Linalg/TransformOps/LinalgMatchOps.cpp.inc"
>();
-
- addDialectDataInitializer<transform::PatternRegistry>(
- [&](transform::PatternRegistry ®istry) {
- 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
diff --git a/mlir/lib/Dialect/Linalg/TransformOps/LinalgTransformOps.cpp b/mlir/lib/Dialect/Linalg/TransformOps/LinalgTransformOps.cpp
index 133ce91bbcb84..4934a60f578b5 100644
--- a/mlir/lib/Dialect/Linalg/TransformOps/LinalgTransformOps.cpp
+++ b/mlir/lib/Dialect/Linalg/TransformOps/LinalgTransformOps.cpp
@@ -141,9 +141,34 @@ static DiagnosedSilenceableFailure unpackSingleIndexResultPayloadOperations(
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) {
diff --git a/mlir/lib/Dialect/MemRef/TransformOps/MemRefTransformOps.cpp b/mlir/lib/Dialect/MemRef/TransformOps/MemRefTransformOps.cpp
index 7b636137e20ac..0f84fe498d911 100644
--- a/mlir/lib/Dialect/MemRef/TransformOps/MemRefTransformOps.cpp
+++ b/mlir/lib/Dialect/MemRef/TransformOps/MemRefTransformOps.cpp
@@ -16,7 +16,6 @@
#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"
@@ -27,6 +26,35 @@ 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
//===----------------------------------------------------------------------===//
@@ -72,31 +100,6 @@ DiagnosedSilenceableFailure transform::MemRefMultiBufferOp::apply(
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
//===----------------------------------------------------------------------===//
@@ -162,23 +165,6 @@ class MemRefTransformDialectExtension
#define GET_OP_LIST
#include "mlir/Dialect/MemRef/TransformOps/MemRefTransformOps.cpp.inc"
>();
-
- addDialectDataInitializer<transform::PatternRegistry>(
- [&](transform::PatternRegistry ®istry) {
- 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
diff --git a/mlir/lib/Dialect/SCF/TransformOps/SCFTransformOps.cpp b/mlir/lib/Dialect/SCF/TransformOps/SCFTransformOps.cpp
index 4c64c5420bf9e..edd1d328d57bc 100644
--- a/mlir/lib/Dialect/SCF/TransformOps/SCFTransformOps.cpp
+++ b/mlir/lib/Dialect/SCF/TransformOps/SCFTransformOps.cpp
@@ -22,6 +22,15 @@
using namespace mlir;
using namespace mlir::affine;
+//===----------------------------------------------------------------------===//
+// Apply...PatternsOp
+//===----------------------------------------------------------------------===//
+
+void transform::ApplyForLoopCanonicalizationPatternsOp::populatePatterns(
+ RewritePatternSet &patterns) {
+ scf::populateSCFForLoopCanonicalizationPatterns(patterns);
+}
+
//===----------------------------------------------------------------------===//
// GetParentForOp
//===----------------------------------------------------------------------===//
@@ -309,13 +318,6 @@ class SCFTransformDialectExtension
#define GET_OP_LIST
#include "mlir/Dialect/SCF/TransformOps/SCFTransformOps.cpp.inc"
>();
-
- addDialectDataInitializer<transform::PatternRegistry>(
- [&](transform::PatternRegistry ®istry) {
- registry.registerPatterns(
- "scf.for_loop_canonicalization",
- scf::populateSCFForLoopCanonicalizationPatterns);
- });
}
};
} // namespace
diff --git a/mlir/lib/Dialect/Tensor/TransformOps/TensorTransformOps.cpp b/mlir/lib/Dialect/Tensor/TransformOps/TensorTransformOps.cpp
index 1ab7883b9be5f..18517cfa93655 100644
--- a/mlir/lib/Dialect/Tensor/TransformOps/TensorTransformOps.cpp
+++ b/mlir/lib/Dialect/Tensor/TransformOps/TensorTransformOps.cpp
@@ -79,6 +79,40 @@ void tensor::registerFindPayloadReplacementOpInterfaceExternalModels(
});
}
+//===----------------------------------------------------------------------===//
+// 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
//===----------------------------------------------------------------------===//
@@ -144,26 +178,6 @@ class TensorTransformDialectExtension
#define GET_OP_LIST
#include "mlir/Dialect/Tensor/TransformOps/TensorTransformOps.cpp.inc"
>();
-
- addDialectDataInitializer<transform::PatternRegistry>(
- [&](transform::PatternRegistry ®istry) {
- 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
diff --git a/mlir/lib/Dialect/Transform/IR/TransformOps.cpp b/mlir/lib/Dialect/Transform/IR/TransformOps.cpp
index 212deb137077f..e2029560e67ee 100644
--- a/mlir/lib/Dialect/Transform/IR/TransformOps.cpp
+++ b/mlir/lib/Dialect/Transform/IR/TransformOps.cpp
@@ -32,8 +32,6 @@
using namespace mlir;
-MLIR_DEFINE_EXPLICIT_TYPE_ID(mlir::transform::PatternRegistry)
-
static ParseResult parseSequenceOpOperands(
OpAsmParser &parser, std::optional<OpAsmParser::UnresolvedOperand> &root,
Type &rootType,
@@ -217,37 +215,6 @@ void transform::ErrorCheckingTrackingListener::notifyPayloadReplacementNotFound(
++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
//===----------------------------------------------------------------------===//
@@ -440,11 +407,6 @@ transform::ApplyPatternsOp::applyToOne(Operation *target,
// Gather all specified patterns.
MLIRContext *ctx = target->getContext();
RewritePatternSet patterns(ctx);
- const auto ®istry = 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(
@@ -495,17 +457,6 @@ transform::ApplyPatternsOp::applyToOne(Operation *target,
}
LogicalResult transform::ApplyPatternsOp::verify() {
- const auto ®istry = 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)) {
diff --git a/mlir/lib/Dialect/Vector/TransformOps/VectorTransformOps.cpp b/mlir/lib/Dialect/Vector/TransformOps/VectorTransformOps.cpp
index 44caaec2d1910..7e097942c8189 100644
--- a/mlir/lib/Dialect/Vector/TransformOps/VectorTransformOps.cpp
+++ b/mlir/lib/Dialect/Vector/TransformOps/VectorTransformOps.cpp
@@ -24,7 +24,7 @@ using namespace mlir::vector;
using namespace mlir::transform;
//===----------------------------------------------------------------------===//
-// ApplyRankReducingSubviewPatternsOp
+// Apply...PatternsOp
//===----------------------------------------------------------------------===//
void transform::ApplyRankReducingSubviewPatternsOp::populatePatterns(
@@ -32,29 +32,17 @@ void transform::ApplyRankReducingSubviewPatternsOp::populatePatterns(
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());
@@ -63,37 +51,23 @@ void transform::LowerContractionOp::populatePatterns(
/*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());
@@ -101,38 +75,33 @@ void transform::LowerMultiReductionOp::populatePatterns(
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(
@@ -148,22 +117,15 @@ void transform::LowerTransposeOp::populatePatterns(
}
}
-//===----------------------------------------------------------------------===//
-// 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())
@@ -189,34 +151,6 @@ class VectorTransformDialectExtension
#define GET_OP_LIST
#include "mlir/Dialect/Vector/TransformOps/VectorTransformOps.cpp.inc"
>();
-
- addDialectDataInitializer<transform::PatternRegistry>(
- [&](transform::PatternRegistry ®istry) {
- 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
diff --git a/mlir/test/Dialect/LLVM/transform-e2e.mlir b/mlir/test/Dialect/LLVM/transform-e2e.mlir
index 3fda2748c058a..049e893b53fa5 100644
--- a/mlir/test/Dialect/LLVM/transform-e2e.mlir
+++ b/mlir/test/Dialect/LLVM/transform-e2e.mlir
@@ -27,35 +27,35 @@ transform.sequence failures(propagate) {
// 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
}
diff --git a/mlir/test/Dialect/Linalg/transform-op-matmul-to-outerproduct.mlir b/mlir/test/Dialect/Linalg/transform-op-matmul-to-outerproduct.mlir
index f3140a9bda8a6..6ebcdd463988b 100644
--- a/mlir/test/Dialect/Linalg/transform-op-matmul-to-outerproduct.mlir
+++ b/mlir/test/Dialect/Linalg/transform-op-matmul-to-outerproduct.mlir
@@ -32,7 +32,7 @@ transform.sequence failures(propagate) {
%0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!transform.any_op) -> !transform.any_op
%1 = get_closest_isolated_parent %0 : (!transform.any_op) -> !transform.any_op
%2 = transform.structured.vectorize %1 : (!transform.any_op) -> !transform.any_op
- transform.apply_patterns [] to %2 {
+ transform.apply_patterns to %2 {
transform.apply_patterns.vector.lower_contraction lowering_strategy = "outerproduct"
} : !transform.any_op
}
diff --git a/mlir/test/Dialect/MemRef/extract-address-computations.mlir b/mlir/test/Dialect/MemRef/extract-address-computations.mlir
index 5064f60d2f7b0..40393ff24e4a5 100644
--- a/mlir/test/Dialect/MemRef/extract-address-computations.mlir
+++ b/mlir/test/Dialect/MemRef/extract-address-computations.mlir
@@ -24,9 +24,11 @@ func.func @test_load(%base : memref<2x16x16xf32>, %offset : index) -> 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
}
// -----
@@ -50,7 +52,9 @@ func.func @test_load_nontemporal(%base : memref<2x16x16xf32>, %offset : index) -
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
}
// -----
@@ -79,7 +83,9 @@ func.func @test_store(%base : memref<2x16x16xf32>, %offset : index) -> () {
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
}
// -----
@@ -105,7 +111,9 @@ func.func @test_store_nontemporal(%base : memref<2x16x16xf32>, %offset : index)
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
}
// -----
@@ -159,7 +167,9 @@ func.func @testWithLoop(%base : memref<?x?x?xf32, strided<[?,?,?], offset: ?>>)
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
}
// -----
@@ -197,7 +207,9 @@ func.func @test_ldmatrix(%base : memref<4x32x32xf16, 3>,
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
}
// -----
@@ -231,7 +243,9 @@ func.func @test_ldmatrix(%base : memref<?x?x?xf16, 3>,
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
}
// -----
@@ -266,7 +280,9 @@ func.func @test_transfer_read_op(%base : memref<?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
}
// -----
@@ -294,7 +310,9 @@ func.func @test_transfer_read_op_with_tensor(%base : 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
}
// -----
@@ -328,7 +346,9 @@ func.func @test_transfer_write_op(%base : memref<?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
}
// -----
@@ -363,8 +383,11 @@ func.func @test_transfer_write_op_with_strides(%base : memref<?x?x?xf16, strided
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.
@@ -389,5 +412,8 @@ func.func @test_transfer_write_op_with_tensor(%base : 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
}
+
diff --git a/mlir/test/Dialect/Transform/ops-invalid.mlir b/mlir/test/Dialect/Transform/ops-invalid.mlir
index 427b1b174acbb..8aa76140ddcb4 100644
--- a/mlir/test/Dialect/Transform/ops-invalid.mlir
+++ b/mlir/test/Dialect/Transform/ops-invalid.mlir
@@ -675,25 +675,10 @@ module attributes { transform.with_named_sequence } {
// -----
-transform.sequence failures(propagate) {
-^bb0(%arg0: !transform.any_op):
- // expected-error @below {{patterns not registered: transform.invalid_pattern_identifier}}
- transform.apply_patterns ["transform.invalid_pattern_identifier"] to %arg0 {} : !transform.any_op
-}
-
-// -----
-
-transform.sequence failures(propagate) {
-^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
diff --git a/mlir/test/Dialect/Transform/test-pattern-application.mlir b/mlir/test/Dialect/Transform/test-pattern-application.mlir
index 55bb083eb9833..ca277abcc9a8b 100644
--- a/mlir/test/Dialect/Transform/test-pattern-application.mlir
+++ b/mlir/test/Dialect/Transform/test-pattern-application.mlir
@@ -15,29 +15,7 @@ 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.
@@ -60,7 +38,9 @@ transform.sequence failures(propagate) {
%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
@@ -84,7 +64,9 @@ transform.sequence failures(propagate) {
%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
}
// -----
@@ -104,7 +86,9 @@ 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
}
@@ -120,7 +104,9 @@ func.func @patterns_apply_only_to_target_body() {
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
}
// -----
@@ -142,7 +128,9 @@ transform.sequence failures(propagate) {
%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
}
@@ -162,7 +150,7 @@ 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
diff --git a/mlir/test/Dialect/Vector/transform-vector.mlir b/mlir/test/Dialect/Vector/transform-vector.mlir
index 7bdbaf9514a32..15f423f854d93 100644
--- a/mlir/test/Dialect/Vector/transform-vector.mlir
+++ b/mlir/test/Dialect/Vector/transform-vector.mlir
@@ -30,35 +30,35 @@ transform.sequence failures(propagate) {
// 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
}
diff --git a/mlir/test/Dialect/Vector/vector-broadcast-lowering-transforms.mlir b/mlir/test/Dialect/Vector/vector-broadcast-lowering-transforms.mlir
index c71b6d6d8a7a7..2d3c88d751192 100644
--- a/mlir/test/Dialect/Vector/vector-broadcast-lowering-transforms.mlir
+++ b/mlir/test/Dialect/Vector/vector-broadcast-lowering-transforms.mlir
@@ -167,7 +167,7 @@ transform.sequence failures(propagate) {
%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
}
diff --git a/mlir/test/Dialect/Vector/vector-contract-matvec-transforms.mlir b/mlir/test/Dialect/Vector/vector-contract-matvec-transforms.mlir
index f897ed0944ef6..d8365f53b43a1 100644
--- a/mlir/test/Dialect/Vector/vector-contract-matvec-transforms.mlir
+++ b/mlir/test/Dialect/Vector/vector-contract-matvec-transforms.mlir
@@ -210,7 +210,7 @@ func.func @redpar_vecmattrans2x2(%arg0: memref<vector<2x2xf32>>, %arg1: memref<v
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
}
diff --git a/mlir/test/Dialect/Vector/vector-contract-to-dot-transforms.mlir b/mlir/test/Dialect/Vector/vector-contract-to-dot-transforms.mlir
index 913b5b3751b1a..28896a157a9dc 100644
--- a/mlir/test/Dialect/Vector/vector-contract-to-dot-transforms.mlir
+++ b/mlir/test/Dialect/Vector/vector-contract-to-dot-transforms.mlir
@@ -300,7 +300,7 @@ transform.sequence failures(propagate) {
%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
}
diff --git a/mlir/test/Dialect/Vector/vector-contract-to-matrix-intrinsics-transforms.mlir b/mlir/test/Dialect/Vector/vector-contract-to-matrix-intrinsics-transforms.mlir
index 2de95789cbddc..e5582e3c3cd78 100644
--- a/mlir/test/Dialect/Vector/vector-contract-to-matrix-intrinsics-transforms.mlir
+++ b/mlir/test/Dialect/Vector/vector-contract-to-matrix-intrinsics-transforms.mlir
@@ -48,11 +48,11 @@ transform.sequence failures(propagate) {
%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
}
diff --git a/mlir/test/Dialect/Vector/vector-contract-to-outerproduct-transforms.mlir b/mlir/test/Dialect/Vector/vector-contract-to-outerproduct-transforms.mlir
index a53064a99d5c6..aee4cf6b1379b 100644
--- a/mlir/test/Dialect/Vector/vector-contract-to-outerproduct-transforms.mlir
+++ b/mlir/test/Dialect/Vector/vector-contract-to-outerproduct-transforms.mlir
@@ -347,7 +347,7 @@ transform.sequence failures(propagate) {
%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
}
diff --git a/mlir/test/Dialect/Vector/vector-contract-to-parallel-arith-transforms.mlir b/mlir/test/Dialect/Vector/vector-contract-to-parallel-arith-transforms.mlir
index 5a7502fbac355..147f3ae921991 100644
--- a/mlir/test/Dialect/Vector/vector-contract-to-parallel-arith-transforms.mlir
+++ b/mlir/test/Dialect/Vector/vector-contract-to-parallel-arith-transforms.mlir
@@ -56,7 +56,7 @@ transform.sequence failures(propagate) {
%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
}
diff --git a/mlir/test/Dialect/Vector/vector-mask-lowering-transforms.mlir b/mlir/test/Dialect/Vector/vector-mask-lowering-transforms.mlir
index d0d53a8356c8a..d425e9c9ab1a6 100644
--- a/mlir/test/Dialect/Vector/vector-mask-lowering-transforms.mlir
+++ b/mlir/test/Dialect/Vector/vector-mask-lowering-transforms.mlir
@@ -96,7 +96,7 @@ transform.sequence failures(propagate) {
%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
}
@@ -127,7 +127,7 @@ transform.sequence failures(propagate) {
%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
}
diff --git a/mlir/test/Dialect/Vector/vector-multi-reduction-lowering.mlir b/mlir/test/Dialect/Vector/vector-multi-reduction-lowering.mlir
index 884cf2840a3fd..5a67c266ee26b 100644
--- a/mlir/test/Dialect/Vector/vector-multi-reduction-lowering.mlir
+++ b/mlir/test/Dialect/Vector/vector-multi-reduction-lowering.mlir
@@ -267,7 +267,7 @@ func.func @vector_multi_reduction_parallel_middle(%arg0: vector<3x4x5xf32>, %acc
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
}
diff --git a/mlir/test/Dialect/Vector/vector-multi-reduction-outer-lowering.mlir b/mlir/test/Dialect/Vector/vector-multi-reduction-outer-lowering.mlir
index eee55e834a9d1..9f0b3a18d67c9 100644
--- a/mlir/test/Dialect/Vector/vector-multi-reduction-outer-lowering.mlir
+++ b/mlir/test/Dialect/Vector/vector-multi-reduction-outer-lowering.mlir
@@ -190,7 +190,7 @@ func.func @vector_multi_reduction_to_scalar(%arg0: vector<2x3xf32>, %acc: 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.lower_multi_reduction lowering_strategy = "innerparallel"
} : !transform.any_op
}
diff --git a/mlir/test/Dialect/Vector/vector-outerproduct-lowering-transforms.mlir b/mlir/test/Dialect/Vector/vector-outerproduct-lowering-transforms.mlir
index 907bdb45a3a68..f6dd80302b4b9 100644
--- a/mlir/test/Dialect/Vector/vector-outerproduct-lowering-transforms.mlir
+++ b/mlir/test/Dialect/Vector/vector-outerproduct-lowering-transforms.mlir
@@ -140,11 +140,11 @@ transform.sequence failures(propagate) {
%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
}
diff --git a/mlir/test/Dialect/Vector/vector-shape-cast-lowering-transforms.mlir b/mlir/test/Dialect/Vector/vector-shape-cast-lowering-transforms.mlir
index 92b3e80b266b0..9ad0bbcfe8b12 100644
--- a/mlir/test/Dialect/Vector/vector-shape-cast-lowering-transforms.mlir
+++ b/mlir/test/Dialect/Vector/vector-shape-cast-lowering-transforms.mlir
@@ -154,7 +154,7 @@ transform.sequence failures(propagate) {
%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
}
diff --git a/mlir/test/Dialect/Vector/vector-transfer-drop-unit-dims-patterns.mlir b/mlir/test/Dialect/Vector/vector-transfer-drop-unit-dims-patterns.mlir
index 749afa238d69a..d64f8889d6820 100644
--- a/mlir/test/Dialect/Vector/vector-transfer-drop-unit-dims-patterns.mlir
+++ b/mlir/test/Dialect/Vector/vector-transfer-drop-unit-dims-patterns.mlir
@@ -17,7 +17,7 @@ func.func @transfer_read_rank_reducing(
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
}
@@ -39,7 +39,7 @@ func.func @transfer_write_rank_reducing(%arg : memref<1x1x3x2xi8, strided<[6, 6,
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
}
@@ -63,7 +63,7 @@ func.func @transfer_read_and_vector_rank_reducing(
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
}
@@ -87,7 +87,7 @@ func.func @transfer_write_and_vector_rank_reducing(
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
}
@@ -111,7 +111,7 @@ func.func @transfer_read_and_vector_rank_reducing_to_0d(
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
}
@@ -135,7 +135,7 @@ func.func @transfer_write_and_vector_rank_reducing_to_0d(
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
}
diff --git a/mlir/test/Dialect/Vector/vector-transfer-full-partial-split-copy-transform.mlir b/mlir/test/Dialect/Vector/vector-transfer-full-partial-split-copy-transform.mlir
index 424c53c93575d..ecb29266246ed 100644
--- a/mlir/test/Dialect/Vector/vector-transfer-full-partial-split-copy-transform.mlir
+++ b/mlir/test/Dialect/Vector/vector-transfer-full-partial-split-copy-transform.mlir
@@ -108,7 +108,7 @@ func.func @split_vector_transfer_read_strided_2d(
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
}
@@ -169,7 +169,7 @@ func.func @split_vector_transfer_write_2d(%V: vector<4x8xf32>, %A: memref<?x8xf3
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
}
@@ -237,7 +237,7 @@ func.func @split_vector_transfer_write_strided_2d(
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
}
diff --git a/mlir/test/Dialect/Vector/vector-transfer-full-partial-split.mlir b/mlir/test/Dialect/Vector/vector-transfer-full-partial-split.mlir
index e7adcfaed41c2..a035c22d75eef 100644
--- a/mlir/test/Dialect/Vector/vector-transfer-full-partial-split.mlir
+++ b/mlir/test/Dialect/Vector/vector-transfer-full-partial-split.mlir
@@ -103,7 +103,7 @@ func.func @split_vector_transfer_read_strided_2d(
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
}
@@ -161,7 +161,7 @@ func.func @split_vector_transfer_write_2d(%V: vector<4x8xf32>, %A: memref<?x8xf3
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
}
@@ -223,7 +223,7 @@ func.func @split_vector_transfer_write_strided_2d(
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
}
@@ -265,7 +265,7 @@ func.func @transfer_read_within_scf_for(%A : memref<?x?xf32>, %lb : index, %ub :
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
}
diff --git a/mlir/test/Dialect/Vector/vector-transfer-to-vector-load-store.mlir b/mlir/test/Dialect/Vector/vector-transfer-to-vector-load-store.mlir
index e136219a2e03c..ad9bec92173b1 100644
--- a/mlir/test/Dialect/Vector/vector-transfer-to-vector-load-store.mlir
+++ b/mlir/test/Dialect/Vector/vector-transfer-to-vector-load-store.mlir
@@ -240,7 +240,7 @@ func.func @transfer_broadcasting_complex(%mem : memref<10x20x30x8x8xf32>, %i : i
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
@@ -362,7 +362,7 @@ func.func @transfer_write_broadcast_unit_dim(
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
diff --git a/mlir/test/Dialect/Vector/vector-transpose-lowering.mlir b/mlir/test/Dialect/Vector/vector-transpose-lowering.mlir
index 0865f97a9fd4d..f8b36055f1002 100644
--- a/mlir/test/Dialect/Vector/vector-transpose-lowering.mlir
+++ b/mlir/test/Dialect/Vector/vector-transpose-lowering.mlir
@@ -76,7 +76,7 @@ func.func @transpose1023_1x1x8x8xf32(%arg0: vector<1x1x8x8xf32>) -> vector<1x1x8
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
}
@@ -99,7 +99,7 @@ func.func @transpose(%arg0: vector<2x4xf32>) -> 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
}
@@ -118,7 +118,7 @@ func.func @transpose(%arg0: vector<2x4xf32>) -> 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
}
@@ -605,7 +605,7 @@ func.func @transpose210_1x8x8xf32(%arg0: vector<1x8x8xf32>) -> 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
}
@@ -683,7 +683,7 @@ func.func @transpose_shuffle16x16xf32(%arg0: vector<16x16xf32>) -> vector<16x16x
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
}
@@ -762,7 +762,7 @@ func.func @transpose021_shuffle16x16xf32(%arg0: vector<1x16x16xf32>) -> vector<1
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
}
diff --git a/mlir/test/Integration/Dialect/Vector/CPU/test-shuffle16x16.mlir b/mlir/test/Integration/Dialect/Vector/CPU/test-shuffle16x16.mlir
index 798a224285f00..74d47de05b706 100644
--- a/mlir/test/Integration/Dialect/Vector/CPU/test-shuffle16x16.mlir
+++ b/mlir/test/Integration/Dialect/Vector/CPU/test-shuffle16x16.mlir
@@ -31,7 +31,7 @@ func.func @entry() {
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
}
diff --git a/mlir/test/lib/Dialect/Transform/TestTransformDialectExtension.cpp b/mlir/test/lib/Dialect/Transform/TestTransformDialectExtension.cpp
index 9243723947b6d..fc6f323a3b557 100644
--- a/mlir/test/lib/Dialect/Transform/TestTransformDialectExtension.cpp
+++ b/mlir/test/lib/Dialect/Transform/TestTransformDialectExtension.cpp
@@ -784,10 +784,6 @@ void mlir::test::ApplyTestPatternsOp::populatePatterns(
}
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.
@@ -825,11 +821,6 @@ class TestTransformDialectExtension
constraints.try_emplace("verbose_constraint", verboseConstraint);
hooks.mergeInPDLMatchHooks(std::move(constraints));
});
-
- addDialectDataInitializer<transform::PatternRegistry>(
- [&](transform::PatternRegistry ®istry) {
- registry.registerPatterns("transform.test", populateTestPatterns);
- });
}
};
} // namespace
More information about the Mlir-commits
mailing list