[Mlir-commits] [mlir] e1c5cbc - [mlir][Linalg] Put a proper type on transform.structured.match op
Nicolas Vasilache
llvmlistbot at llvm.org
Thu Jan 26 08:53:02 PST 2023
Author: Nicolas Vasilache
Date: 2023-01-26T08:51:34-08:00
New Revision: e1c5cbc09c6fe6ba396638ff96a10796225c97e8
URL: https://github.com/llvm/llvm-project/commit/e1c5cbc09c6fe6ba396638ff96a10796225c97e8
DIFF: https://github.com/llvm/llvm-project/commit/e1c5cbc09c6fe6ba396638ff96a10796225c97e8.diff
LOG: [mlir][Linalg] Put a proper type on transform.structured.match op
This allows much better verification messages in consuming ops that properly declare
`TransformHandleTypeInterface` on their operands.
Downstream tests can be updated with a command resembling:
```
git grep -l "structured\.match" mlir/test | xargs -i sed -i {} -e "s/\(structured.match.*\)/\1 : (\!pdl.operation) -> \!pdl.operation/g"
```
Differential Revision: https://reviews.llvm.org/D142643
Added:
Modified:
mlir/include/mlir/Dialect/Linalg/TransformOps/LinalgTransformOps.td
mlir/test/Dialect/Affine/transform-op-simplify-bounded-affine-ops.mlir
mlir/test/Dialect/Bufferization/Transforms/transform-ops.mlir
mlir/test/Dialect/GPU/transform-gpu-failing.mlir
mlir/test/Dialect/GPU/transform-gpu.mlir
mlir/test/Dialect/LLVM/transform-e2e.mlir
mlir/test/Dialect/Linalg/generalize-tensor-pack.mlir
mlir/test/Dialect/Linalg/generalize-tensor-unpack.mlir
mlir/test/Dialect/Linalg/multisize-tiling-full.mlir
mlir/test/Dialect/Linalg/promote.mlir
mlir/test/Dialect/Linalg/promotion_options.mlir
mlir/test/Dialect/Linalg/tile-conv.mlir
mlir/test/Dialect/Linalg/tile-indexed.mlir
mlir/test/Dialect/Linalg/tile-tensors.mlir
mlir/test/Dialect/Linalg/tile-to-foreach-thread.mlir
mlir/test/Dialect/Linalg/transform-op-decompose.mlir
mlir/test/Dialect/Linalg/transform-op-fuse-into-containing.mlir
mlir/test/Dialect/Linalg/transform-op-fuse.mlir
mlir/test/Dialect/Linalg/transform-op-generalize.mlir
mlir/test/Dialect/Linalg/transform-op-interchange.mlir
mlir/test/Dialect/Linalg/transform-op-match.mlir
mlir/test/Dialect/Linalg/transform-op-multitile-sizes.mlir
mlir/test/Dialect/Linalg/transform-op-pack.mlir
mlir/test/Dialect/Linalg/transform-op-pad.mlir
mlir/test/Dialect/Linalg/transform-op-replace.mlir
mlir/test/Dialect/Linalg/transform-op-scalarize.mlir
mlir/test/Dialect/Linalg/transform-op-split-reduction-by-scaling.mlir
mlir/test/Dialect/Linalg/transform-op-split-reduction.mlir
mlir/test/Dialect/Linalg/transform-op-split.mlir
mlir/test/Dialect/Linalg/transform-op-tile.mlir
mlir/test/Dialect/Linalg/transform-op-vectorize.mlir
mlir/test/Dialect/Linalg/transform-patterns.mlir
mlir/test/Dialect/Linalg/transform-promotion.mlir
mlir/test/Dialect/Linalg/transform-tile-and-fuse.mlir
mlir/test/Dialect/Linalg/transform-tile-reduction.mlir
mlir/test/Dialect/Linalg/vectorization.mlir
mlir/test/Dialect/MemRef/transform-ops.mlir
mlir/test/Dialect/SCF/transform-op-coalesce.mlir
mlir/test/Dialect/SCF/transform-ops-invalid.mlir
mlir/test/Dialect/SCF/transform-ops.mlir
mlir/test/Dialect/Tensor/tiling.mlir
mlir/test/Dialect/Transform/test-interpreter.mlir
mlir/test/Dialect/Vector/transform-vector.mlir
mlir/test/Integration/Dialect/Linalg/CPU/test-conv-1d-call.mlir
mlir/test/Integration/Dialect/Linalg/CPU/test-conv-1d-nwc-wcf-call.mlir
mlir/test/Integration/Dialect/Linalg/CPU/test-conv-2d-call.mlir
mlir/test/Integration/Dialect/Linalg/CPU/test-conv-2d-nhwc-hwcf-call.mlir
mlir/test/Integration/Dialect/Linalg/CPU/test-conv-3d-call.mlir
mlir/test/Integration/Dialect/Linalg/CPU/test-conv-3d-ndhwc-dhwcf-call.mlir
mlir/test/Integration/Dialect/Linalg/CPU/test-tensor-matmul.mlir
Removed:
################################################################################
diff --git a/mlir/include/mlir/Dialect/Linalg/TransformOps/LinalgTransformOps.td b/mlir/include/mlir/Dialect/Linalg/TransformOps/LinalgTransformOps.td
index 804ce5b370eb1..30506fadb70f7 100644
--- a/mlir/include/mlir/Dialect/Linalg/TransformOps/LinalgTransformOps.td
+++ b/mlir/include/mlir/Dialect/Linalg/TransformOps/LinalgTransformOps.td
@@ -260,13 +260,13 @@ def MatchOp : Op<Transform_Dialect, "structured.match",
it is a navigation op.
}];
- let arguments = (ins PDL_Operation:$target,
+ let arguments = (ins TransformHandleTypeInterface:$target,
OptionalAttr<StrArrayAttr>:$ops,
OptionalAttr<MatchInterfaceEnum>:$interface,
OptionalAttr<DictionaryAttr>:$op_attrs,
OptionalAttr<TypeAttr>:$filter_result_type);
// TODO: variadic results when needed.
- let results = (outs PDL_Operation:$results);
+ let results = (outs TransformHandleTypeInterface:$results);
let builders = [
OpBuilder<(ins "Value":$target, "ArrayRef<StringRef>":$opNames)>
@@ -278,6 +278,7 @@ def MatchOp : Op<Transform_Dialect, "structured.match",
(`attributes` $op_attrs^)?
(`filter_result_type` `=` $filter_result_type^)?
`in` $target attr-dict
+ `:` functional-type($target, results)
}];
}
diff --git a/mlir/test/Dialect/Affine/transform-op-simplify-bounded-affine-ops.mlir b/mlir/test/Dialect/Affine/transform-op-simplify-bounded-affine-ops.mlir
index ec9559c081bcf..2f384ea3a251c 100644
--- a/mlir/test/Dialect/Affine/transform-op-simplify-bounded-affine-ops.mlir
+++ b/mlir/test/Dialect/Affine/transform-op-simplify-bounded-affine-ops.mlir
@@ -15,8 +15,8 @@ func.func @simplify_min_max() -> (index, index) {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["affine.min", "affine.max"]} in %arg1
- %1 = transform.structured.match ops{["test.some_op"]} in %arg1
+ %0 = transform.structured.match ops{["affine.min", "affine.max"]} in %arg1 : (!pdl.operation) -> !pdl.operation
+ %1 = transform.structured.match ops{["test.some_op"]} in %arg1 : (!pdl.operation) -> !pdl.operation
transform.affine.simplify_bounded_affine_ops %0 with [%1] within [0] and [20]
}
@@ -35,9 +35,9 @@ func.func @simplify_min_sequence() -> index {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["affine.min"]} in %arg1
- %1 = transform.structured.match ops{["test.workgroup_id"]} in %arg1
- %2 = transform.structured.match ops{["test.thread_id"]} in %arg1
+ %0 = transform.structured.match ops{["affine.min"]} in %arg1 : (!pdl.operation) -> !pdl.operation
+ %1 = transform.structured.match ops{["test.workgroup_id"]} in %arg1 : (!pdl.operation) -> !pdl.operation
+ %2 = transform.structured.match ops{["test.thread_id"]} in %arg1 : (!pdl.operation) -> !pdl.operation
transform.affine.simplify_bounded_affine_ops %0 with [%1, %2] within [0, 0] and [31, 31]
}
@@ -45,7 +45,7 @@ transform.sequence failures(propagate) {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["affine.min"]} in %arg1
+ %0 = transform.structured.match ops{["affine.min"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error at +1 {{incorrect number of lower bounds, expected 0 but found 1}}
transform.affine.simplify_bounded_affine_ops %0 with [] within [0] and []
}
@@ -54,7 +54,7 @@ transform.sequence failures(propagate) {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["affine.min"]} in %arg1
+ %0 = transform.structured.match ops{["affine.min"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error at +1 {{incorrect number of upper bounds, expected 0 but found 1}}
transform.affine.simplify_bounded_affine_ops %0 with [] within [] and [5]
}
diff --git a/mlir/test/Dialect/Bufferization/Transforms/transform-ops.mlir b/mlir/test/Dialect/Bufferization/Transforms/transform-ops.mlir
index ec537b6828e31..b71a1c0cfcfa2 100644
--- a/mlir/test/Dialect/Bufferization/Transforms/transform-ops.mlir
+++ b/mlir/test/Dialect/Bufferization/Transforms/transform-ops.mlir
@@ -4,7 +4,7 @@
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["func.func"]} in %arg1
+ %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!pdl.operation) -> !pdl.operation
transform.bufferization.one_shot_bufferize %0
{target_is_module = false}
}
@@ -33,7 +33,7 @@ func.func @test_function(%A : tensor<?xf32>, %v : vector<4xf32>) -> (tensor<?xf3
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["func.func"]} in %arg1
+ %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!pdl.operation) -> !pdl.operation
transform.bufferization.one_shot_bufferize %0
{target_is_module = false, test_analysis_only = true}
}
@@ -56,7 +56,7 @@ func.func @test_function_analysis(%A : tensor<?xf32>, %v : vector<4xf32>) -> (te
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["func.func"]} in %arg1
+ %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error @+1 {{bufferization failed}}
transform.bufferization.one_shot_bufferize %0 {target_is_module = false}
}
@@ -123,7 +123,7 @@ func.func @matmul(%A: tensor<12x9xf32>, %B: tensor<9x6xf32>, %C: tensor<12x6xf32
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.empty"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.empty"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = transform.cast %0 : !pdl.operation to !transform.op<"tensor.empty">
transform.bufferization.empty_tensor_to_alloc_tensor %1 : (!transform.op<"tensor.empty">) -> !transform.op<"bufferization.alloc_tensor">
}
diff --git a/mlir/test/Dialect/GPU/transform-gpu-failing.mlir b/mlir/test/Dialect/GPU/transform-gpu-failing.mlir
index bc414a31c6083..e48b48a060b85 100644
--- a/mlir/test/Dialect/GPU/transform-gpu-failing.mlir
+++ b/mlir/test/Dialect/GPU/transform-gpu-failing.mlir
@@ -6,7 +6,7 @@ func.func @map_nested_foreach_to_threads_not_gpu_launch() -> () {
}
transform.sequence failures(propagate) {
^bb0(%arg0: !pdl.operation):
- %funcop = transform.structured.match ops{["tensor.empty"]} in %arg0
+ %funcop = transform.structured.match ops{["tensor.empty"]} in %arg0 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{Given target is not gpu.launch}}
%1 = transform.gpu.map_nested_foreach_to_threads %funcop
}
@@ -46,7 +46,7 @@ func.func @map_nested_foreach_to_threads_excessive_threads(%x: memref<2 x 32 x f
}
transform.sequence failures(propagate) {
^bb1(%arg0: !pdl.operation):
- %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0
+ %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{Trying to launch a GPU kernel with gridDim = (1, 1, 1) blockDim = (1200, 9, 1). It is larger than the limits.}}
// expected-note @below {{"blockDim" is very large}}
transform.gpu.map_nested_foreach_to_threads %funcop { blockDim = [1200, 9, 1] }
@@ -88,7 +88,7 @@ func.func @map_nested_foreach_to_threads_fewer_threads(%x: memref<2 x 32 x f32>,
transform.sequence failures(propagate) {
^bb1(%arg0: !pdl.operation):
- %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0
+ %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{The requested GPU threads are fewer than the number of loop trip counts. Try to tile scf.foreach_thread before mapping or set small blockDim.}}
transform.gpu.map_nested_foreach_to_threads %funcop { blockDim = [128, 4, 1] }
}
@@ -114,7 +114,7 @@ func.func @map_nested_foreach_to_threads_dynamic_trip_count(%x: memref<2 x 32 x
transform.sequence failures(propagate) {
^bb1(%arg0: !pdl.operation):
- %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0
+ %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{unsupported dynamic blockdim size}}
transform.gpu.map_nested_foreach_to_threads %funcop { blockDim = [128, 4, 1] }
}
@@ -134,9 +134,9 @@ func.func @map_nested_foreach_to_threads_not_buffer(%x: tensor<32x32xf32>, %y: t
transform.sequence failures(propagate) {
^bb1(%arg0: !pdl.operation):
- %matmul = transform.structured.match ops{["linalg.matmul"]} in %arg0
+ %matmul = transform.structured.match ops{["linalg.matmul"]} in %arg0 : (!pdl.operation) -> !pdl.operation
%foreach, %tiled = transform.structured.tile_to_foreach_thread_op %matmul num_threads [10, 20, 30] (mapping = [ #gpu.thread<y>, #gpu.thread<x>, #gpu.thread<z> ] )
- %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0
+ %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{only bufferized scf.foreach_thread lowers to gpu.thread_id}}
transform.gpu.map_nested_foreach_to_threads %funcop { blockDim = [128, 4, 1] }
}
@@ -151,7 +151,7 @@ func.func @map_foreach_to_blocks_not_gpu_launch() -> () {
}
transform.sequence failures(propagate) {
^bb0(%arg0: !pdl.operation):
- %funcop = transform.structured.match ops{["tensor.empty"]} in %arg0
+ %funcop = transform.structured.match ops{["tensor.empty"]} in %arg0 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{Given target is not gpu.launch}}
%1 = transform.gpu.map_foreach_to_blocks %funcop
}
@@ -188,7 +188,7 @@ func.func @map_foreach_to_blocks_not_unique(%x: memref<2 x 32 x f32>, %y: memref
transform.sequence failures(propagate) {
^bb0(%arg0: !pdl.operation):
- %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0
+ %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{could not find a unique topLevel scf.foreach_thread}}
%1 = transform.gpu.map_foreach_to_blocks %funcop
}
@@ -221,7 +221,7 @@ func.func @map_foreach_to_blocks_large_loop(%x: memref<2 x 32 x f32>, %y: memref
transform.sequence failures(propagate) {
^bb0(%arg0: !pdl.operation):
- %funcop = transform.structured.match ops{["func.func"]} in %arg0
+ %funcop = transform.structured.match ops{["func.func"]} in %arg0 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{could not find a unique topLevel scf.foreach_thread}}
%1 = transform.gpu.map_foreach_to_blocks %funcop { generate_gpu_launch }
}
@@ -242,7 +242,7 @@ func.func @map_foreach_to_blocks_large_loop(%x: memref<2 x 32 x f32>, %y: memref
transform.sequence failures(propagate) {
^bb0(%arg0: !pdl.operation):
- %funcop = transform.structured.match ops{["func.func"]} in %arg0
+ %funcop = transform.structured.match ops{["func.func"]} in %arg0 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{Trying to launch a GPU kernel with gridDim = (65535, 65535, 1) blockDim = (1, 1, 1). It is larger than the limits.}}
%1 = transform.gpu.map_foreach_to_blocks %funcop { generate_gpu_launch }
}
@@ -269,7 +269,7 @@ func.func @saxpy2d_singleloop(%x: !type, %y: !type, %stream : !gpu.async.token)
transform.sequence failures(propagate) {
^bb1(%arg0: !pdl.operation):
- %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0
+ %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{#gpu.thread<x> is duplicated, cannot map
diff erent loops to the same processor}}
transform.gpu.map_nested_foreach_to_threads %funcop { blockDim = [32, 32]}
}
@@ -296,7 +296,7 @@ func.func @saxpy2d_wrong_mapping(%x: !type, %y: !type, %stream : !gpu.async.toke
transform.sequence failures(propagate) {
^bb1(%arg0: !pdl.operation):
- %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0
+ %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{mapping must be one of #gpu.thread<x>, #gpu.thread<y>, #gpu.thread<z>}}
transform.gpu.map_nested_foreach_to_threads %funcop { blockDim = [32, 32]}
}
diff --git a/mlir/test/Dialect/GPU/transform-gpu.mlir b/mlir/test/Dialect/GPU/transform-gpu.mlir
index 1c337fed5cc09..24b474cbd10d6 100644
--- a/mlir/test/Dialect/GPU/transform-gpu.mlir
+++ b/mlir/test/Dialect/GPU/transform-gpu.mlir
@@ -32,7 +32,7 @@ func.func @saxpy2dblock(%x: !type, %y: !type, %t: !type1d, %alpha : f32, %stream
transform.sequence failures(propagate) {
^bb1(%arg0: !pdl.operation):
- %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0
+ %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!pdl.operation) -> !pdl.operation
transform.gpu.map_foreach_to_blocks %funcop { gridDim = [12, 9]}
}
@@ -86,7 +86,7 @@ func.func @saxpy2d(%x: !type, %y: !type, %t: !type1d, %alpha : f32, %stream : !g
transform.sequence failures(propagate) {
^bb1(%arg0: !pdl.operation):
- %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0
+ %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!pdl.operation) -> !pdl.operation
transform.gpu.map_nested_foreach_to_threads %funcop { blockDim = [12, 9] }
}
@@ -125,7 +125,7 @@ func.func @saxpy4d(%x: !type4d, %y: !type4d, %alpha : f32) -> !type4d {
transform.sequence failures(propagate) {
^bb1(%arg0: !pdl.operation):
- %funcop = transform.structured.match ops{["func.func"]} in %arg0
+ %funcop = transform.structured.match ops{["func.func"]} in %arg0 : (!pdl.operation) -> !pdl.operation
%gpuLaunch = transform.gpu.map_foreach_to_blocks %funcop { generate_gpu_launch }
transform.gpu.map_nested_foreach_to_threads %gpuLaunch { blockDim = [32, 4, 1] }
}
@@ -159,7 +159,7 @@ func.func @saxpy2d_no_barrier(%x: !type, %y: !type, %t: !type1d, %alpha : f32, %
transform.sequence failures(propagate) {
^bb1(%arg0: !pdl.operation):
- %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0
+ %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!pdl.operation) -> !pdl.operation
transform.gpu.map_nested_foreach_to_threads %funcop { blockDim = [12, 9, 1], syncAfterDistribute = false }
}
@@ -191,7 +191,7 @@ func.func @saxpy2d_singleloop(%x: !type, %y: !type, %stream : !gpu.async.token)
transform.sequence failures(propagate) {
^bb1(%arg0: !pdl.operation):
- %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0
+ %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!pdl.operation) -> !pdl.operation
transform.gpu.map_nested_foreach_to_threads %funcop { blockDim = [32]}
}
@@ -227,6 +227,6 @@ func.func @saxpy3d_fold_id_z(%x: !type, %y: !type, %t: !type1d, %alpha : f32, %s
transform.sequence failures(propagate) {
^bb1(%arg0: !pdl.operation):
- %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0
+ %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!pdl.operation) -> !pdl.operation
transform.gpu.map_nested_foreach_to_threads %funcop { blockDim = [12, 9, 1], syncAfterDistribute = false }
}
diff --git a/mlir/test/Dialect/LLVM/transform-e2e.mlir b/mlir/test/Dialect/LLVM/transform-e2e.mlir
index 9fae730a2bd16..d091e9d18d1b5 100644
--- a/mlir/test/Dialect/LLVM/transform-e2e.mlir
+++ b/mlir/test/Dialect/LLVM/transform-e2e.mlir
@@ -14,12 +14,12 @@ func.func @matmul_tensors(
transform.sequence failures(propagate) {
^bb1(%module_op: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %module_op
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %module_op : (!pdl.operation) -> !pdl.operation
%1, %loops:3 = transform.structured.tile %0 [2, 2, 2] : (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation, !pdl.operation)
%2 = get_closest_isolated_parent %1 : (!pdl.operation) -> !pdl.operation
transform.structured.vectorize %2
transform.bufferization.one_shot_bufferize layout{IdentityLayoutMap} %module_op
{bufferize_function_boundaries = true}
- %func = transform.structured.match ops{["func.func"]} in %module_op
+ %func = transform.structured.match ops{["func.func"]} in %module_op : (!pdl.operation) -> !pdl.operation
transform.vector.lower_vectors %func multireduction_lowering = "innerreduction"
}
diff --git a/mlir/test/Dialect/Linalg/generalize-tensor-pack.mlir b/mlir/test/Dialect/Linalg/generalize-tensor-pack.mlir
index 0db5d60b72a10..8d3bd2b805569 100644
--- a/mlir/test/Dialect/Linalg/generalize-tensor-pack.mlir
+++ b/mlir/test/Dialect/Linalg/generalize-tensor-pack.mlir
@@ -93,7 +93,7 @@ func.func @KCRS_to_KCRSsr(%arg0: tensor<1x1x128x64xf32>, %arg1: tensor<1x1x4x8x8
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.pack"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.pack"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:4 = transform.structured.tile_to_scf_for %0 [1, 1, 1, 1]
}
@@ -122,7 +122,7 @@ func.func @pad_and_pack(%arg0: tensor<13x15xf32>, %arg1: tensor<2x8x8x2xf32>, %a
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.pack"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.pack"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:2 = transform.structured.tile_to_scf_for %0 [1, 1]
}
@@ -161,6 +161,6 @@ func.func @KC_to_CKkc(%arg0: tensor<128x256xf32>, %arg1: tensor<32x4x32x8xf32>)
// CHECK-TRANS-SAME: [%[[C]], %[[K]], 0, 0] [1, 1, 32, 8] [1, 1, 1, 1] : tensor<1x1x32x8xf32> into tensor<32x4x32x8xf32>
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.pack"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.pack"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:2 = transform.structured.tile_to_scf_for %0 [1, 1]
}
diff --git a/mlir/test/Dialect/Linalg/generalize-tensor-unpack.mlir b/mlir/test/Dialect/Linalg/generalize-tensor-unpack.mlir
index f8e223d8de528..a03161f9648fe 100644
--- a/mlir/test/Dialect/Linalg/generalize-tensor-unpack.mlir
+++ b/mlir/test/Dialect/Linalg/generalize-tensor-unpack.mlir
@@ -67,7 +67,7 @@ func.func @KCRSsr_to_KCRS(%arg0: tensor<1x1x4x8x8x32xf32>, %arg1: tensor<1x1x128
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.unpack"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.unpack"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:4 = transform.structured.tile_to_scf_for %0 [1, 1, 32, 8]
}
// CHECK-TRANS-DAG: #[[MAP0:.+]] = affine_map<(d0) -> (d0 floordiv 32)>
@@ -129,7 +129,7 @@ func.func @unpack_and_extract_slice(%arg0: tensor<2x8x8x2xf32>, %arg1: tensor<13
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.unpack"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.unpack"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:2 = transform.structured.tile_to_scf_for %0 [8, 2]
}
@@ -163,6 +163,6 @@ func.func @CKkc_to_KC(%arg0: tensor<32x4x32x8xf32>, %arg1: tensor<128x256xf32>)
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.unpack"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.unpack"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:2 = transform.structured.tile_to_scf_for %0 [32, 8]
}
diff --git a/mlir/test/Dialect/Linalg/multisize-tiling-full.mlir b/mlir/test/Dialect/Linalg/multisize-tiling-full.mlir
index 2e6167651dc95..a5ab963ad8906 100644
--- a/mlir/test/Dialect/Linalg/multisize-tiling-full.mlir
+++ b/mlir/test/Dialect/Linalg/multisize-tiling-full.mlir
@@ -4,7 +4,7 @@
// This implements a 2D multisize tiling with target sizes [3, 10].
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1:3 = transform.structured.multitile_sizes %0 { dimension = 0, target_size = 3} : (!pdl.operation) -> !pdl.operation
%t:3 = transform.structured.multitile_sizes %0 { dimension = 1, target_size = 10} : (!pdl.operation) -> !pdl.operation
%2:2 = transform.structured.split %0 after %1#2 { dimension = 0 } : !pdl.operation, !pdl.operation
@@ -104,7 +104,7 @@ func.func @two_d(%arg0: tensor<10x34xf32>,
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1:3 = transform.structured.multitile_sizes %0 { dimension = 0, target_size = 3} : (!pdl.operation) -> !transform.param<i64>
%t:3 = transform.structured.multitile_sizes %0 { dimension = 1, target_size = 10} : (!pdl.operation) -> !transform.param<i64>
%2:2 = transform.structured.split %0 after %1#2 { dimension = 0 } : !pdl.operation, !transform.param<i64>
diff --git a/mlir/test/Dialect/Linalg/promote.mlir b/mlir/test/Dialect/Linalg/promote.mlir
index 1a384d0d68529..085c1b7ae714b 100644
--- a/mlir/test/Dialect/Linalg/promote.mlir
+++ b/mlir/test/Dialect/Linalg/promote.mlir
@@ -68,7 +68,7 @@ func.func @matmul_f32(%A: memref<?xi8>, %M: index, %N: index, %K: index) {
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = transform.structured.promote %0 { use_alloca }
}
@@ -138,7 +138,7 @@ func.func @matmul_f64(%A: memref<?xi8>, %M: index, %N: index, %K: index) {
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = transform.structured.promote %0
}
@@ -184,6 +184,6 @@ func.func @promote_rank_reducing_subviews(%arg0: memref<?x?x?x64xf32, strided<[
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match interface{LinalgOp} in %arg1
+ %0 = transform.structured.match interface{LinalgOp} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = transform.structured.promote %0
}
diff --git a/mlir/test/Dialect/Linalg/promotion_options.mlir b/mlir/test/Dialect/Linalg/promotion_options.mlir
index 70cac58acaf35..f18d448ae9c3a 100644
--- a/mlir/test/Dialect/Linalg/promotion_options.mlir
+++ b/mlir/test/Dialect/Linalg/promotion_options.mlir
@@ -36,7 +36,7 @@ func.func @gemm(%a : memref<?x?xf32>, %b : memref<?x?xf32>, %c : memref<?x?xf32>
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:3 = transform.structured.tile %0 [16, 16, 16] : (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation, !pdl.operation)
%2 = transform.structured.promote %1 { operands_to_promote = [0, 2], force_full_tiles = [false, false], use_full_tiles_by_default }
}
diff --git a/mlir/test/Dialect/Linalg/tile-conv.mlir b/mlir/test/Dialect/Linalg/tile-conv.mlir
index cd126be25dfc2..f7b0b22baa8d0 100644
--- a/mlir/test/Dialect/Linalg/tile-conv.mlir
+++ b/mlir/test/Dialect/Linalg/tile-conv.mlir
@@ -11,7 +11,7 @@ func.func @conv(%arg0 : memref<?x?xf32>, %arg1 : memref<?x?xf32>, %arg2 : memref
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.conv_2d"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.conv_2d"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loop:2 = transform.structured.tile %0 [2, 3] : (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation)
}
diff --git a/mlir/test/Dialect/Linalg/tile-indexed.mlir b/mlir/test/Dialect/Linalg/tile-indexed.mlir
index 3e6e0d5c5c686..68fdc3bc0cc40 100644
--- a/mlir/test/Dialect/Linalg/tile-indexed.mlir
+++ b/mlir/test/Dialect/Linalg/tile-indexed.mlir
@@ -13,7 +13,7 @@ func.func @indexed_vector(%arg0: memref<50xindex>) {
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loop = transform.structured.tile %0 [10] : (!pdl.operation) -> (!pdl.operation, !pdl.operation)
}
@@ -43,7 +43,7 @@ func.func @indexed_matrix(%arg0: memref<50x50xindex>) {
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loop:2 = transform.structured.tile %0 [10, 25] : (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation)
}
diff --git a/mlir/test/Dialect/Linalg/tile-tensors.mlir b/mlir/test/Dialect/Linalg/tile-tensors.mlir
index 484534e740ab3..c192a59087c11 100644
--- a/mlir/test/Dialect/Linalg/tile-tensors.mlir
+++ b/mlir/test/Dialect/Linalg/tile-tensors.mlir
@@ -29,7 +29,7 @@ func.func @matmul_tensors(
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:3 = transform.structured.tile %0 [2, 3, 4] : (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation, !pdl.operation)
}
@@ -60,7 +60,7 @@ func.func @generic_op_tensors(
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:3 = transform.structured.tile %0 [2, 3, 4] : (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation, !pdl.operation)
}
@@ -131,6 +131,6 @@ func.func @fold_extract_slice(
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:3 = transform.structured.tile %0 [2, 3, 4] : (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation, !pdl.operation)
}
diff --git a/mlir/test/Dialect/Linalg/tile-to-foreach-thread.mlir b/mlir/test/Dialect/Linalg/tile-to-foreach-thread.mlir
index 53a7828a81d2b..586d09c51ebac 100644
--- a/mlir/test/Dialect/Linalg/tile-to-foreach-thread.mlir
+++ b/mlir/test/Dialect/Linalg/tile-to-foreach-thread.mlir
@@ -34,7 +34,7 @@ module {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1:2 = transform.structured.tile_to_foreach_thread_op %0 num_threads [10, 20] (mapping = [ #gpu.thread<y>, #gpu.thread<x> ] )
}
}
@@ -76,8 +76,8 @@ func.func @matmul_tile_size_dynamic_dynamic(%A: tensor<?x?xf32>, %B: tensor<?x?x
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
- %sz = transform.structured.match ops{["test.dummy"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
+ %sz = transform.structured.match ops{["test.dummy"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1:2 = transform.structured.tile_to_foreach_thread_op %0 tile_sizes %sz
}
@@ -117,7 +117,7 @@ func.func @matmul_static(%A: tensor<100x200xf32>, %B: tensor<200x300xf32>, %C: t
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1:2 = transform.structured.tile_to_foreach_thread_op %0 num_threads [10, 21]
}
@@ -158,7 +158,7 @@ func.func @matmul_tile_size_dynamic(%A: tensor<?x?xf32>, %B: tensor<?x?xf32>, %C
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1:2 = transform.structured.tile_to_foreach_thread_op %0 tile_sizes [10, 20]
}
@@ -196,7 +196,7 @@ func.func @matmul_tile_size_static(%A: tensor<100x200xf32>, %B: tensor<200x300xf
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1:2 = transform.structured.tile_to_foreach_thread_op %0 tile_sizes [10, 21]
}
@@ -218,7 +218,7 @@ module {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1:2 = transform.structured.tile_to_foreach_thread_op %0 num_threads [2] ( mapping = [#gpu.thread<x>])
}
}
@@ -269,8 +269,8 @@ func.func @matmul_tile_size_dynamic_dynamic(%A: tensor<?x?xf32>, %B: tensor<?x?x
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
- %sz = transform.structured.match ops{["test.dummy"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
+ %sz = transform.structured.match ops{["test.dummy"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1:2 = transform.structured.tile_to_foreach_thread_op %0 tile_sizes [%sz, 20]
}
@@ -324,7 +324,7 @@ transform.sequence failures(propagate) {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%foreach_thread, %tiled_generic = transform.structured.tile_to_foreach_thread_op %0 num_threads [7]
}
@@ -378,7 +378,7 @@ transform.sequence failures(propagate) {
transform.sequence failures(propagate) {
^bb1(%IN_MAT2: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %IN_MAT2
+ %0 = transform.structured.match ops{["linalg.generic"]} in %IN_MAT2 : (!pdl.operation) -> !pdl.operation
%foreach_thread, %tiled_generic = transform.structured.tile_to_foreach_thread_op %0 num_threads [4]
}
diff --git a/mlir/test/Dialect/Linalg/transform-op-decompose.mlir b/mlir/test/Dialect/Linalg/transform-op-decompose.mlir
index 2c873b263bdbe..e023e64c7cc8b 100644
--- a/mlir/test/Dialect/Linalg/transform-op-decompose.mlir
+++ b/mlir/test/Dialect/Linalg/transform-op-decompose.mlir
@@ -184,6 +184,6 @@ func.func @pooling_nchw_max(%input: tensor<?x?x1x?xf32>, %filter: tensor<1x?xf32
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match interface{LinalgOp} in %arg1
+ %0 = transform.structured.match interface{LinalgOp} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = transform.structured.decompose %0
}
diff --git a/mlir/test/Dialect/Linalg/transform-op-fuse-into-containing.mlir b/mlir/test/Dialect/Linalg/transform-op-fuse-into-containing.mlir
index 7424e08f5338c..67ac3b432f8ce 100644
--- a/mlir/test/Dialect/Linalg/transform-op-fuse-into-containing.mlir
+++ b/mlir/test/Dialect/Linalg/transform-op-fuse-into-containing.mlir
@@ -43,8 +43,8 @@ module {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.fill"]} in %arg1
- %1 = transform.structured.match ops{["scf.foreach_thread"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.fill"]} in %arg1 : (!pdl.operation) -> !pdl.operation
+ %1 = transform.structured.match ops{["scf.foreach_thread"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// linalg.fill is tileable. The op is tiled and fused.
transform.structured.fuse_into_containing_op %0 into %1
@@ -86,8 +86,8 @@ module {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.empty"]} in %arg1
- %1 = transform.structured.match ops{["scf.foreach_thread"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.empty"]} in %arg1 : (!pdl.operation) -> !pdl.operation
+ %1 = transform.structured.match ops{["scf.foreach_thread"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// tensor.empty is not tileable. The op is cloned and fused.
transform.structured.fuse_into_containing_op %0 into %1
@@ -132,8 +132,8 @@ module {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.fill"]} in %arg1
- %1 = transform.structured.match ops{["scf.foreach_thread"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.fill"]} in %arg1 : (!pdl.operation) -> !pdl.operation
+ %1 = transform.structured.match ops{["scf.foreach_thread"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// linalg.fill is tileable. The op is tiled and fused.
transform.structured.fuse_into_containing_op %0 into %1
@@ -180,8 +180,8 @@ module {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.fill"]} in %arg1
- %1 = transform.structured.match ops{["scf.foreach_thread"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.fill"]} in %arg1 : (!pdl.operation) -> !pdl.operation
+ %1 = transform.structured.match ops{["scf.foreach_thread"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// linalg.fill is tileable. The op is tiled and fused.
transform.structured.fuse_into_containing_op %0 into %1
@@ -240,8 +240,8 @@ module {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
- %1 = transform.structured.match ops{["scf.foreach_thread"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
+ %1 = transform.structured.match ops{["scf.foreach_thread"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// linalg.generic is tileable. The op is tiled and fused.
transform.structured.fuse_into_containing_op %0 into %1
diff --git a/mlir/test/Dialect/Linalg/transform-op-fuse.mlir b/mlir/test/Dialect/Linalg/transform-op-fuse.mlir
index c50c8f4087e33..66328a18fa5de 100644
--- a/mlir/test/Dialect/Linalg/transform-op-fuse.mlir
+++ b/mlir/test/Dialect/Linalg/transform-op-fuse.mlir
@@ -17,7 +17,7 @@ func.func @fuse_unary(%arg0: tensor<?x?xf32>, %arg1: tensor<?x?xf32>) -> tensor<
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.elemwise_binary"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.elemwise_binary"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:2 = transform.structured.fuse %0 {tile_sizes = [32, 32], tile_interchange = [0, 1]}
}
@@ -44,7 +44,7 @@ func.func @fuse_unary(%arg0: tensor<?x?xf32>, %arg1: tensor<?x?xf32>) -> tensor<
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.elemwise_binary"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.elemwise_binary"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:2 = transform.structured.fuse %0 {tile_sizes = [32, 32], tile_interchange = [0, 1]}
%loop = transform.cast %loops#0 : !pdl.operation to !transform.op<"scf.for">
transform.loop.peel %loop : (!transform.op<"scf.for">) -> !pdl.operation
@@ -87,7 +87,7 @@ func.func @interchange_reduction(%input: tensor<12x7x25xf32>) -> tensor<12x25xf3
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:2 = transform.structured.fuse %0 {tile_sizes = [5, 0, 7], tile_interchange = [0, 2, 1]}
%2, %loops_2 = transform.structured.tile %1 [0, 4] : (!pdl.operation) -> (!pdl.operation, !pdl.operation)
}
@@ -111,6 +111,6 @@ func.func @unpack_elemwise(%arg0: tensor<16x48x8x8xf32>, %arg1: tensor<128x384xf
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.elemwise_unary"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.elemwise_unary"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:2 = transform.structured.fuse %0 {tile_sizes = [16, 32], tile_interchange = [0, 1]}
}
diff --git a/mlir/test/Dialect/Linalg/transform-op-generalize.mlir b/mlir/test/Dialect/Linalg/transform-op-generalize.mlir
index b961494393064..c83ef4c6224a9 100644
--- a/mlir/test/Dialect/Linalg/transform-op-generalize.mlir
+++ b/mlir/test/Dialect/Linalg/transform-op-generalize.mlir
@@ -12,6 +12,6 @@ func.func @generalize_unary(%arg0: tensor<?x?xf32>, %arg1: tensor<?x?xf32>) -> t
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.elemwise_unary"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.elemwise_unary"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = transform.structured.generalize %0
}
diff --git a/mlir/test/Dialect/Linalg/transform-op-interchange.mlir b/mlir/test/Dialect/Linalg/transform-op-interchange.mlir
index 3b480d70f0961..52b636c6bd3f1 100644
--- a/mlir/test/Dialect/Linalg/transform-op-interchange.mlir
+++ b/mlir/test/Dialect/Linalg/transform-op-interchange.mlir
@@ -20,7 +20,7 @@ func.func @interchange_generic(%arg0: tensor<?x?xf32>, %arg1: tensor<?x?xf32>) -
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
transform.structured.interchange %0 iterator_interchange = [1, 0]
}
@@ -34,7 +34,7 @@ func.func @interchange_matmul(%arg0: tensor<?x?xf32>, %arg1: tensor<?x?xf32>, %a
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{transform applied to the wrong op kind}}
transform.structured.interchange %0 iterator_interchange = [1, 0]
}
diff --git a/mlir/test/Dialect/Linalg/transform-op-match.mlir b/mlir/test/Dialect/Linalg/transform-op-match.mlir
index b4ad820d453b7..e41ffec038ed4 100644
--- a/mlir/test/Dialect/Linalg/transform-op-match.mlir
+++ b/mlir/test/Dialect/Linalg/transform-op-match.mlir
@@ -11,11 +11,11 @@ func.func @bar() {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %match_name = transform.structured.match ops{["arith.constant"]} in %arg1
+ %match_name = transform.structured.match ops{["arith.constant"]} in %arg1 : (!pdl.operation) -> !pdl.operation
transform.test_print_remark_at_operand %match_name, "matched op name" : !pdl.operation
transform.test_consume_operand %match_name
- %match_attr = transform.structured.match ops{["arith.constant"]} attributes{my_attr} in %arg1
+ %match_attr = transform.structured.match ops{["arith.constant"]} attributes{my_attr} in %arg1 : (!pdl.operation) -> !pdl.operation
transform.test_print_remark_at_operand %match_attr, "matched attr name" : !pdl.operation
transform.test_consume_operand %match_attr
}
@@ -32,7 +32,7 @@ func.func @by_type() {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
%match_name = transform.structured.match
- ops{["arith.constant"]} filter_result_type = f32 in %arg1
+ ops{["arith.constant"]} filter_result_type = f32 in %arg1 : (!pdl.operation) -> !pdl.operation
transform.test_print_remark_at_operand %match_name, "matched op name" : !pdl.operation
transform.test_consume_operand %match_name
}
@@ -63,7 +63,7 @@ transform.sequence failures(propagate) {
#linalg.iterator_type<parallel>,
#linalg.iterator_type<parallel>,
#linalg.iterator_type<parallel>]}
- in %arg1
+ in %arg1 : (!pdl.operation) -> !pdl.operation
transform.test_print_remark_at_operand %match_attr, "matched complex attr" : !pdl.operation
transform.test_consume_operand %match_attr
@@ -72,7 +72,7 @@ transform.sequence failures(propagate) {
#linalg.iterator_type<parallel>,
#linalg.iterator_type<parallel>,
#linalg.iterator_type<reduction>]}
- in %arg1
+ in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-remark @below {{0}}
transform.test_print_number_of_associated_payload_ir_ops %no_match
}
diff --git a/mlir/test/Dialect/Linalg/transform-op-multitile-sizes.mlir b/mlir/test/Dialect/Linalg/transform-op-multitile-sizes.mlir
index b035d58d18864..d9fa8cbc30078 100644
--- a/mlir/test/Dialect/Linalg/transform-op-multitile-sizes.mlir
+++ b/mlir/test/Dialect/Linalg/transform-op-multitile-sizes.mlir
@@ -4,7 +4,7 @@
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
transform.structured.multitile_sizes %0 { target_size = 3, dimension = 0 } : (!pdl.operation) -> !pdl.operation
}
@@ -28,7 +28,7 @@ func.func @multitile_sizes_static(
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%low_tile, %high_tile, %split_point =
transform.structured.multitile_sizes %0 { target_size = 3, dimension = 0 }
: (!pdl.operation) -> !transform.param<i64>
@@ -55,7 +55,7 @@ func.func @multitile_sizes_static_gen(
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
transform.structured.multitile_sizes %0 { target_size = 3, divisor = 2, dimension = 0 } : (!pdl.operation) -> !pdl.operation
}
@@ -96,7 +96,7 @@ func.func @multitile_sizes_dynamic(
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{cannot compute parametric tile sizes for dynamically shaped payload op}}
transform.structured.multitile_sizes %0 { target_size = 3, divisor = 2, dimension = 0 }
: (!pdl.operation) -> !transform.param<i64>
diff --git a/mlir/test/Dialect/Linalg/transform-op-pack.mlir b/mlir/test/Dialect/Linalg/transform-op-pack.mlir
index b8c569f0a3c17..b2796e3354320 100644
--- a/mlir/test/Dialect/Linalg/transform-op-pack.mlir
+++ b/mlir/test/Dialect/Linalg/transform-op-pack.mlir
@@ -35,7 +35,7 @@ func.func @reduction_2d_static(%t0: tensor<3x7xf16>, %t1: tensor<3xf16>) -> tens
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
transform.structured.pack %0 packed_sizes = [0, 4]
: (!pdl.operation) -> (!transform.op<"linalg.generic">)
}
@@ -77,7 +77,7 @@ func.func @col_reduction_2d_static(%t0: tensor<7x3xf16>, %t1: tensor<3xf16>) ->
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = transform.structured.pack %0 packed_sizes = [4, 0]
: (!pdl.operation) -> (!transform.op<"linalg.generic">)
%pack = transform.get_producer_of_operand %1[0]
@@ -132,7 +132,7 @@ func.func @reduction_2d_dynamic(%t0: tensor<?x?xf16>, %t1: tensor<?xf16>) -> ten
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
transform.structured.pack %0 packed_sizes = [0, 4]
: (!pdl.operation) -> (!transform.op<"linalg.generic">)
}
@@ -178,7 +178,7 @@ func.func @reduction_2d_dynamic(%t0: tensor<?x?xf16>, %t1: tensor<?xf16>) -> ten
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
transform.structured.pack %0 packed_sizes = [3, 4]
: (!pdl.operation) -> (!transform.op<"linalg.generic">)
}
@@ -221,7 +221,7 @@ func.func @matmul(%A: tensor<?x?xf32>, %B: tensor<?x?xf32>, %C: tensor<?x?xf32>)
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// M N K
%1 = transform.structured.pack %0 packed_sizes = [2, 3, 4]
: (!pdl.operation) -> (!transform.op<"linalg.generic">)
@@ -269,7 +269,7 @@ func.func @conv_2d_nchw_fchw(%i: tensor<14x512x28x28xf32>, %f: tensor<1024x512x1
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match interface{LinalgOp} in %arg1
+ %0 = transform.structured.match interface{LinalgOp} in %arg1 : (!pdl.operation) -> !pdl.operation
// N F H W C KH KW
%1 = transform.structured.pack %0 packed_sizes = [0, 4, 0, 0, 8, 0, 0]
: (!pdl.operation) -> (!transform.op<"linalg.generic">)
@@ -310,7 +310,7 @@ func.func @conv_2d_nhwc_hwcf(%input: tensor<?x1x?x?xf32>, %filter: tensor<1x?x?x
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match interface{LinalgOp} in %arg1
+ %0 = transform.structured.match interface{LinalgOp} in %arg1 : (!pdl.operation) -> !pdl.operation
// N H W F KH KW C
%1 = transform.structured.pack %0 packed_sizes = [0, 0, 0, 4, 0, 0, 6]
: (!pdl.operation) -> (!transform.op<"linalg.generic">)
@@ -356,8 +356,8 @@ func.func @matmul_dynamic_pack_size(%A: tensor<?x?xf32>, %B: tensor<?x?xf32>, %C
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
- %sz = transform.structured.match ops{["some_tile_size"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
+ %sz = transform.structured.match ops{["some_tile_size"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = transform.structured.pack %0 packed_sizes = [0, %sz, %sz]
: (!pdl.operation) -> (!transform.op<"linalg.generic">)
}
@@ -373,7 +373,7 @@ func.func @conv_cant_pack(%i: tensor<14x512x28x28xf32>, %f: tensor<1024x512x1x1x
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match interface{LinalgOp} in %arg1
+ %0 = transform.structured.match interface{LinalgOp} in %arg1 : (!pdl.operation) -> !pdl.operation
// N F H W C KH KW
// expected-error @below {{data tiling failed}}
%1 = transform.structured.pack %0 packed_sizes = [0, 0, 4, 0, 0, 0, 0]
@@ -395,7 +395,7 @@ func.func @matmul(%A: tensor<?x?xf32>, %B: tensor<?x?xf32>, %C: tensor<?x?xf32>)
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{requires target to map to exactly 1 LinalgOp (got 2)}}
%1 = transform.structured.pack %0 packed_sizes = [2, 3, 4]
: (!pdl.operation) -> (!transform.op<"linalg.generic">)
@@ -414,7 +414,7 @@ func.func @matmul(%A: tensor<?x?xf32>, %B: tensor<?x?xf32>, %C: tensor<?x?xf32>)
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{requires number of packed sizes match the number of loops (2 vs 3)}}
%1 = transform.structured.pack %0 packed_sizes = [2, 3]
: (!pdl.operation) -> (!transform.op<"linalg.generic">)
@@ -431,8 +431,8 @@ func.func @no_single_packing_op(%source: tensor<128x256xf32>, %dest: tensor<4x16
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.pack"]} in %arg1
- %1 = transform.structured.match ops{["tensor.unpack"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.pack"]} in %arg1 : (!pdl.operation) -> !pdl.operation
+ %1 = transform.structured.match ops{["tensor.unpack"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{requires target to map to exactly 1 packing op and 1 packed op (got 2 and 1)}}
transform.structured.pack_transpose %0 with_compute_op(%1)
inner_perm = [0]
@@ -450,8 +450,8 @@ func.func @no_single_pack_unpack(%source: tensor<128x256xf32>, %dest: tensor<4x1
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["arith.constant"]} in %arg1
- %1 = transform.structured.match ops{["tensor.empty"]} in %arg1
+ %0 = transform.structured.match ops{["arith.constant"]} in %arg1 : (!pdl.operation) -> !pdl.operation
+ %1 = transform.structured.match ops{["tensor.empty"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{requires target to map to a tensor.pack or tensor.unpack}}
transform.structured.pack_transpose %0 with_compute_op(%1)
inner_perm = [0]
@@ -469,8 +469,8 @@ func.func @no_linalg_target(%source: tensor<128x256xf32>, %dest: tensor<4x16x32x
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.pack"]} in %arg1
- %1 = transform.structured.match ops{["arith.constant"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.pack"]} in %arg1 : (!pdl.operation) -> !pdl.operation
+ %1 = transform.structured.match ops{["arith.constant"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{requires a LinalgOp target}}
transform.structured.pack_transpose %0 with_compute_op(%1)
inner_perm = [0]
@@ -490,8 +490,8 @@ func.func @no_single_use_by_linalg(%source: tensor<128x256xf32>, %dest: tensor<4
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.pack"]} in %arg1
- %1 = transform.structured.match ops{["linalg.fill"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.pack"]} in %arg1 : (!pdl.operation) -> !pdl.operation
+ %1 = transform.structured.match ops{["linalg.fill"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{not a single use by the LinalgOp target}}
transform.structured.pack_transpose %0 with_compute_op(%1)
inner_perm = [0]
@@ -512,8 +512,8 @@ func.func @not_produced_by_linalg(%source: tensor<128x256xf32>, %dest: tensor<4x
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.unpack"]} in %arg1
- %1 = transform.structured.match ops{["linalg.fill"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.unpack"]} in %arg1 : (!pdl.operation) -> !pdl.operation
+ %1 = transform.structured.match ops{["linalg.fill"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{not produced by the LinalgOp target}}
transform.structured.pack_transpose %0 with_compute_op(%1)
inner_perm = [0]
@@ -533,8 +533,8 @@ func.func @no_matching_pack(%source: tensor<16xf32>) {
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.unpack"]} in %arg1
- %1 = transform.structured.match ops{["linalg.fill"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.unpack"]} in %arg1 : (!pdl.operation) -> !pdl.operation
+ %1 = transform.structured.match ops{["linalg.fill"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{could not find matching pack op}}
transform.structured.pack_transpose %0 with_compute_op(%1)
inner_perm = [0]
@@ -554,7 +554,7 @@ func.func @invalid_outer_perm(%A: tensor<?x?xf32>, %B: tensor<?x?xf32>, %C: tens
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = transform.structured.pack %0 packed_sizes = [2, 3, 4]
: (!pdl.operation) -> (!transform.op<"linalg.generic">)
@@ -580,7 +580,7 @@ func.func @invalid_inner_perm(%A: tensor<?x?xf32>, %B: tensor<?x?xf32>, %C: tens
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = transform.structured.pack %0 packed_sizes = [2, 3, 4]
: (!pdl.operation) -> (!transform.op<"linalg.generic">)
diff --git a/mlir/test/Dialect/Linalg/transform-op-pad.mlir b/mlir/test/Dialect/Linalg/transform-op-pad.mlir
index e391f05c5880e..dcd800300df9d 100644
--- a/mlir/test/Dialect/Linalg/transform-op-pad.mlir
+++ b/mlir/test/Dialect/Linalg/transform-op-pad.mlir
@@ -33,7 +33,7 @@ func.func @static_sizes_output_divisible(%arg0: tensor<24x12xf32>,
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = transform.structured.pad %0 {padding_values=[0.0 : f32, 0.0 : f32, 0.0 : f32], padding_dimensions=[0, 1, 2], pack_paddings=[1, 1, 0]}
}
@@ -49,7 +49,7 @@ func.func @pad(%arg0: tensor<24x12xf32>,
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{op expects a padding value of type 'f32', got 0 : i32}}
%1 = transform.structured.pad %0 {padding_values=[0: i32, 0.0 : f32, 0.0 : f32], padding_dimensions=[0, 1, 2], pack_paddings=[1, 1, 0]}
}
@@ -66,7 +66,7 @@ func.func @pad(%arg0: tensor<24x12xf32>,
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{expects a padding that parses to 'f32', got "foo"}}
%1 = transform.structured.pad %0 {padding_values=["foo", 0.0 : f32, 0.0 : f32], padding_dimensions=[0, 1, 2], pack_paddings=[1, 1, 0]}
}
@@ -84,7 +84,7 @@ func.func @pad(%arg0: tensor<24x12xf32>,
transform.sequence failures(suppress) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// This error is silenceable and is not reported by this transform
// {{transform.structured.pad failed to apply}}
%1 = transform.structured.pad %0 {padding_values=[0.0 : f32, 0.0 : f32, 0.0 : f32], padding_dimensions=[0, 1, 2], pack_paddings=[1, 1, 0]}
diff --git a/mlir/test/Dialect/Linalg/transform-op-replace.mlir b/mlir/test/Dialect/Linalg/transform-op-replace.mlir
index f8ed7712b68e0..3d30a2cfe14f5 100644
--- a/mlir/test/Dialect/Linalg/transform-op-replace.mlir
+++ b/mlir/test/Dialect/Linalg/transform-op-replace.mlir
@@ -10,7 +10,7 @@ func.func @bar() {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["func.func"]} in %arg1
+ %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!pdl.operation) -> !pdl.operation
transform.structured.replace %0 {
func.func @foo() {
"dummy_op"() : () -> ()
@@ -26,7 +26,7 @@ func.func @bar(%arg0: i1) {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["another_op"]} in %arg1
+ %0 = transform.structured.match ops{["another_op"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error @+1 {{expected target without operands}}
transform.structured.replace %0 {
"dummy_op"() : () -> ()
@@ -41,7 +41,7 @@ func.func @bar() {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["another_op"]} in %arg1
+ %0 = transform.structured.match ops{["another_op"]} in %arg1 : (!pdl.operation) -> !pdl.operation
transform.structured.replace %0 {
^bb0(%a: i1):
// expected-error @+1 {{expected replacement without operands}}
diff --git a/mlir/test/Dialect/Linalg/transform-op-scalarize.mlir b/mlir/test/Dialect/Linalg/transform-op-scalarize.mlir
index e46f19d34ae2f..272d0b50333a1 100644
--- a/mlir/test/Dialect/Linalg/transform-op-scalarize.mlir
+++ b/mlir/test/Dialect/Linalg/transform-op-scalarize.mlir
@@ -20,7 +20,7 @@ func.func @scalarize(%arg0: tensor<24x12xf32>,
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops = transform.structured.tile %0 [10, 0, 0] : (!pdl.operation) -> (!pdl.operation, !pdl.operation)
%2 = transform.structured.scalarize %1
}
diff --git a/mlir/test/Dialect/Linalg/transform-op-split-reduction-by-scaling.mlir b/mlir/test/Dialect/Linalg/transform-op-split-reduction-by-scaling.mlir
index 6631bcf2597ac..974a2b763b065 100644
--- a/mlir/test/Dialect/Linalg/transform-op-split-reduction-by-scaling.mlir
+++ b/mlir/test/Dialect/Linalg/transform-op-split-reduction-by-scaling.mlir
@@ -20,7 +20,7 @@ func.func @matmul_split(%A : tensor<?x256xf32>, %B: tensor<256x32xf32>, %C: tens
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1:4 = transform.structured.split_reduction %0
{ split_factor = 4, insert_split_dimension = 2, use_scaling_algorithm, use_alloc}
}
diff --git a/mlir/test/Dialect/Linalg/transform-op-split-reduction.mlir b/mlir/test/Dialect/Linalg/transform-op-split-reduction.mlir
index 7b4d99d126431..e5615d54ca7d8 100644
--- a/mlir/test/Dialect/Linalg/transform-op-split-reduction.mlir
+++ b/mlir/test/Dialect/Linalg/transform-op-split-reduction.mlir
@@ -33,7 +33,7 @@ func.func @matmul_split(%A : tensor<16x256xf32>, %B: tensor<256x32xf32>, %C: ten
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1:4 = transform.structured.split_reduction %0 { split_factor = 4, insert_split_dimension = 2}
}
@@ -81,7 +81,7 @@ func.func @generic_split_1d(%arg0: tensor<32xf32>, %arg1: tensor<f32>, %out: ten
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1:4 = transform.structured.split_reduction %0 { split_factor = 4, insert_split_dimension = 0}
}
@@ -132,7 +132,7 @@ func.func @generic_split_3d(%input: tensor<32x2xf32>, %input_2: tensor<5x32xf32>
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1:4 = transform.structured.split_reduction %0 { split_factor = 4, insert_split_dimension = 2}
}
@@ -171,7 +171,7 @@ func.func @matmul_split(%A : tensor<16x256xf32>, %B: tensor<256x32xf32>, %C: ten
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1:4 = transform.structured.split_reduction %0 { split_factor = 4, insert_split_dimension = 2, inner_parallel}
}
@@ -219,7 +219,7 @@ func.func @generic_split_1d(%arg0: tensor<32xf32>, %arg1: tensor<f32>, %out: ten
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1:4 = transform.structured.split_reduction %0 { split_factor = 4, insert_split_dimension = 0, inner_parallel}
}
@@ -270,6 +270,6 @@ func.func @generic_split_3d(%input: tensor<32x2xf32>, %input_2: tensor<5x32xf32>
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1:4 = transform.structured.split_reduction %0 { split_factor = 4, insert_split_dimension = 2, inner_parallel}
}
diff --git a/mlir/test/Dialect/Linalg/transform-op-split.mlir b/mlir/test/Dialect/Linalg/transform-op-split.mlir
index 6313e77fc9dde..7860881e1b517 100644
--- a/mlir/test/Dialect/Linalg/transform-op-split.mlir
+++ b/mlir/test/Dialect/Linalg/transform-op-split.mlir
@@ -2,7 +2,7 @@
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1:2 = transform.structured.split %0 after 42 { dimension = 0 } : !pdl.operation
}
@@ -50,7 +50,7 @@ func.func @one_d_static(%arg0: tensor<100xf32>, %arg1: tensor<100xf32>) -> tenso
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1:2 = transform.structured.split %0 after 42 { dimension = 0 } : !pdl.operation
}
@@ -83,8 +83,8 @@ func.func @one_d_static_overflow(%arg0: tensor<10xf32>, %arg1: tensor<10xf32>) -
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
- %1 = transform.structured.match ops{["func.call"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
+ %1 = transform.structured.match ops{["func.call"]} in %arg1 : (!pdl.operation) -> !pdl.operation
transform.structured.split %0 after %1 { dimension = 0 } : !pdl.operation, !pdl.operation
}
@@ -131,7 +131,7 @@ func.func @dynamic(%arg0: tensor<100xf32>, %arg1: tensor<100xf32>) -> tensor<100
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1:2 = transform.structured.split %0 after 4 { dimension = 0 } : !pdl.operation
%2:2 = transform.structured.split %1#1 after 16 { dimension = 1 } : !pdl.operation
}
@@ -196,8 +196,8 @@ transform.sequence failures(propagate) {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
- %1 = transform.structured.match ops{["func.call"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
+ %1 = transform.structured.match ops{["func.call"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{expected dynamic split point handle to point to a single-result index-typed op}}
transform.structured.split %0 after %1 { dimension = 0 } : !pdl.operation, !pdl.operation
}
@@ -222,8 +222,8 @@ func.func @dynamic(%arg0: tensor<100xf32>, %arg1: tensor<100xf32>) -> tensor<100
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
- %1 = transform.structured.match ops{["func.call"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
+ %1 = transform.structured.match ops{["func.call"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{expected the dynamic split point handle to point to as many operations (0) as the target handle (1)}}
transform.structured.split %0 after %1 { dimension = 0 } : !pdl.operation, !pdl.operation
}
@@ -246,7 +246,7 @@ func.func @dynamic(%arg0: tensor<100xf32>, %arg1: tensor<100xf32>) -> tensor<100
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["func.return"]} in %arg1
+ %0 = transform.structured.match ops{["func.return"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{only applies to structured ops}}
transform.structured.split %0 after 16 { dimension = 1 } : !pdl.operation
}
@@ -260,7 +260,7 @@ func.func @noop(%arg0: tensor<100xf32>, %arg1: tensor<100xf32>) -> tensor<100xf3
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{dimension 1 does not exist in target op}}
transform.structured.split %0 after 16 { dimension = 1 } : !pdl.operation
}
@@ -282,7 +282,7 @@ func.func @one_d_static(%arg0: tensor<100xf32>, %arg1: tensor<100xf32>) -> tenso
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{splitting does not produce the second part for a subset of targets}}
// expected-note @below {{expected splitting to produce the second part of all or none of the targets}}
%1:2 = transform.structured.split %0 after 142 { dimension = 0 } : !pdl.operation
diff --git a/mlir/test/Dialect/Linalg/transform-op-tile.mlir b/mlir/test/Dialect/Linalg/transform-op-tile.mlir
index 10517050bf01d..df48c2138fbd8 100644
--- a/mlir/test/Dialect/Linalg/transform-op-tile.mlir
+++ b/mlir/test/Dialect/Linalg/transform-op-tile.mlir
@@ -2,7 +2,7 @@
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:3 = transform.structured.tile %0 [4, 4, 4] : (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation, !pdl.operation)
}
@@ -38,8 +38,8 @@ func.func @tile_linalg_matmul(
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
- %1 = transform.structured.match ops{["func.call"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
+ %1 = transform.structured.match ops{["func.call"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%2, %loops:3 = transform.structured.tile %0 [%1, %1, 4] : (!pdl.operation, !pdl.operation, !pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation, !pdl.operation)
}
@@ -78,7 +78,7 @@ func.func @tile_linalg_matmul_dynamic(
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-note @below {{for this parameter}}
%1 = transform.test_produce_integer_param_with_type i64 : !transform.param<i64>
// expected-error @below {{expected as many parameter values (0) as target ops (2)}}
@@ -103,9 +103,9 @@ func.func @tile_linalg_matmul(
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-note @below {{for this handle}}
- %1 = transform.structured.match ops{["arith.constant"]} in %arg1
+ %1 = transform.structured.match ops{["arith.constant"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{expected as many dynamic size-producing operations (0) as target ops (2)}}
transform.structured.tile %0 [%1, %1, 1]
: (!pdl.operation, !pdl.operation, !pdl.operation)
diff --git a/mlir/test/Dialect/Linalg/transform-op-vectorize.mlir b/mlir/test/Dialect/Linalg/transform-op-vectorize.mlir
index 5d34f952a204a..155b0785d2ec7 100644
--- a/mlir/test/Dialect/Linalg/transform-op-vectorize.mlir
+++ b/mlir/test/Dialect/Linalg/transform-op-vectorize.mlir
@@ -18,7 +18,7 @@ func.func @vectorize_matmul(%arg0: tensor<24x12xf32>,
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1
}
@@ -64,7 +64,7 @@ func.func @vectorize_keep_pad(
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1
}
@@ -112,7 +112,7 @@ func.func @vectorize_pad(
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1 {vectorize_padding}
}
@@ -129,7 +129,7 @@ func.func @vectorize(%arg0: tensor<24x12xf32>,
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{op requires isolated-from-above targets}}
%2 = transform.structured.vectorize %0
}
diff --git a/mlir/test/Dialect/Linalg/transform-patterns.mlir b/mlir/test/Dialect/Linalg/transform-patterns.mlir
index c34e94a0b511d..e2063e181480b 100644
--- a/mlir/test/Dialect/Linalg/transform-patterns.mlir
+++ b/mlir/test/Dialect/Linalg/transform-patterns.mlir
@@ -13,7 +13,7 @@ func.func @dot(%x: memref<?xf32, strided<[1], offset: ?>>,
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.dot"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.dot"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loop = transform.structured.tile %0 [8000] : (!pdl.operation) -> (!pdl.operation, !pdl.operation)
}
@@ -37,7 +37,7 @@ func.func @matvec(%A: memref<?x?xf32, strided<[?, 1], offset: ?>>,
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matvec"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matvec"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:2 = transform.structured.tile %0 [5, 6] : (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation)
}
@@ -64,7 +64,7 @@ func.func @matmul(%A: memref<?x?xf32, strided<[?, 1], offset: ?>>,
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:3 = transform.structured.tile %0 [2000, 3000, 4000] : (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation, !pdl.operation)
%2, %loops_2:3 = transform.structured.tile %1 [200, 300, 400] : (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation, !pdl.operation)
%3, %loops_3:3 = transform.structured.tile %2 [20, 30, 40] : (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation, !pdl.operation)
@@ -137,7 +137,7 @@ func.func @permute_generic(%A: memref<?x?xf32, strided<[?, 1], offset: ?>>,
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
transform.structured.interchange %0 iterator_interchange = [1, 2, 0]
}
@@ -163,7 +163,7 @@ func.func @matvec_perm(%A: memref<?x?xf32, strided<[?, 1], offset: ?>>,
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matvec"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matvec"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:2 = transform.structured.tile %0 [5, 6] {interchange = [1, 0]} : (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation)
}
@@ -190,7 +190,7 @@ func.func @matmul_perm(%A: memref<?x?xf32, strided<[?, 1], offset: ?>>,
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:3 = transform.structured.tile %0 [2000, 3000, 4000] {interchange = [1, 2, 0]} : (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation, !pdl.operation)
%2, %loops_2:3 = transform.structured.tile %1 [200, 300, 400] {interchange = [1, 0, 2]} : (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation, !pdl.operation)
%3, %loops_3:3 = transform.structured.tile %2 [20, 30, 40] : (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation, !pdl.operation)
diff --git a/mlir/test/Dialect/Linalg/transform-promotion.mlir b/mlir/test/Dialect/Linalg/transform-promotion.mlir
index ab57b36fbc799..b29a6cbdc28cd 100644
--- a/mlir/test/Dialect/Linalg/transform-promotion.mlir
+++ b/mlir/test/Dialect/Linalg/transform-promotion.mlir
@@ -60,7 +60,7 @@ func.func @promote_subview_matmul(%arg0: memref<?x?xf32, strided<[?, 1], offset:
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = transform.structured.promote %0 { operands_to_promote = [0, 1, 2], use_full_tiles_by_default }
}
@@ -122,7 +122,7 @@ transform.with_pdl_patterns {
^bb0(%arg0: !pdl.operation):
sequence %arg0 : !pdl.operation failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = transform.structured.promote %0 { operands_to_promote = [0], use_full_tiles_by_default }
}
}
@@ -155,7 +155,7 @@ transform.with_pdl_patterns {
^bb0(%arg0: !pdl.operation):
sequence %arg0 : !pdl.operation failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.fill"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.fill"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = transform.structured.promote %0 { operands_to_promote = [1], use_full_tile_buffers = [false, true], alignment = 32}
}
}
@@ -189,7 +189,7 @@ transform.with_pdl_patterns {
^bb0(%arg0: !pdl.operation):
sequence %arg0 : !pdl.operation failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.fill"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.fill"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = transform.structured.promote %0 { operands_to_promote = [1], use_full_tile_buffers = [false, true], alignment = 32}
}
}
diff --git a/mlir/test/Dialect/Linalg/transform-tile-and-fuse.mlir b/mlir/test/Dialect/Linalg/transform-tile-and-fuse.mlir
index d95abcbbfcc79..140ffe437a20d 100644
--- a/mlir/test/Dialect/Linalg/transform-tile-and-fuse.mlir
+++ b/mlir/test/Dialect/Linalg/transform-tile-and-fuse.mlir
@@ -43,8 +43,8 @@ module {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
// Find the root and all producers.
- %root = transform.structured.match attributes{"__root__"} in %arg1
- %producers = transform.structured.match attributes{"__producer__"} in %arg1
+ %root = transform.structured.match attributes{"__root__"} in %arg1 : (!pdl.operation) -> !pdl.operation
+ %producers = transform.structured.match attributes{"__producer__"} in %arg1 : (!pdl.operation) -> !pdl.operation
// Tile the root.
%foreach_thread_op, %tiled_op = transform.structured.tile_to_foreach_thread_op %root num_threads [10, 20]
@@ -100,8 +100,8 @@ module {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
// Find the root and all producers.
- %root = transform.structured.match attributes{"__root__"} in %arg1
- %producers = transform.structured.match attributes{"__producer__"} in %arg1
+ %root = transform.structured.match attributes{"__root__"} in %arg1 : (!pdl.operation) -> !pdl.operation
+ %producers = transform.structured.match attributes{"__producer__"} in %arg1 : (!pdl.operation) -> !pdl.operation
%reversed_producers = transform.test_reverse_payload_ops %producers
// Tile the root.
diff --git a/mlir/test/Dialect/Linalg/transform-tile-reduction.mlir b/mlir/test/Dialect/Linalg/transform-tile-reduction.mlir
index 1a23bcf887e7b..ec8956cf716cc 100644
--- a/mlir/test/Dialect/Linalg/transform-tile-reduction.mlir
+++ b/mlir/test/Dialect/Linalg/transform-tile-reduction.mlir
@@ -16,7 +16,7 @@ func.func @reduction_tile(%arg0: tensor<?x?xf32>, %out: tensor<?xf32>) -> tensor
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%loop, %1, %2, %3 = transform.structured.tile_reduction_using_scf %0
by tile_sizes = [0, 5]
}
@@ -71,7 +71,7 @@ func.func @reduction_tile_transpose(%arg0: tensor<?x?xf32>, %out: tensor<?xf32>)
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%loop, %1, %2, %3 = transform.structured.tile_reduction_using_scf %0
by tile_sizes = [5, 0]
}
@@ -108,7 +108,7 @@ func.func @reduction_tile_parallel(
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%loop, %1, %2, %3 = transform.structured.tile_reduction_using_foreach_thread %0
by num_threads = [0, 5], tile_sizes = []
}
@@ -161,7 +161,7 @@ func.func @matmul_tile_parallel(
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%loop, %1, %2, %3 = transform.structured.tile_reduction_using_foreach_thread %0
by num_threads = [0, 0, 5], tile_sizes = []
}
@@ -221,7 +221,7 @@ func.func @reduction_tile_parallel_cyclic_dist(
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%loop, %1, %2, %3 = transform.structured.tile_reduction_using_foreach_thread %0
by num_threads = [0, 5], tile_sizes = [0, 3], mapping = [#gpu.thread<x>]
}
@@ -287,7 +287,7 @@ func.func @reduction_tile_parallel_cyclic_dist(
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%loop, %1, %2, %3 = transform.structured.tile_reduction_using_foreach_thread %0
by num_threads = [0, 5], tile_sizes = [0, 3], mapping = [#gpu.thread<x>]
@@ -324,7 +324,7 @@ func.func @reduction_untiled_foreach_thread(
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{could not tile reduction}}
%loop, %1, %2, %3 = transform.structured.tile_reduction_using_foreach_thread %0
by num_threads = [5], tile_sizes = [3], mapping = [#gpu.thread<x>]
diff --git a/mlir/test/Dialect/Linalg/vectorization.mlir b/mlir/test/Dialect/Linalg/vectorization.mlir
index 4f34eccec3a93..171a518447697 100644
--- a/mlir/test/Dialect/Linalg/vectorization.mlir
+++ b/mlir/test/Dialect/Linalg/vectorization.mlir
@@ -12,7 +12,7 @@ func.func @contraction_dot(%A: memref<1584xf32>, %B: memref<1584xf32>, %C: memre
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.dot"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.dot"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1 { disable_multi_reduction_to_contract_patterns }
}
@@ -31,7 +31,7 @@ func.func @contraction_matvec(%A: memref<1584x1584xf32>, %B: memref<1584xf32>, %
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matvec"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matvec"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1 { disable_multi_reduction_to_contract_patterns }
}
@@ -49,7 +49,7 @@ func.func @contraction_matmul(%A: memref<1584x1584xf32>, %B: memref<1584x1584xf3
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1 { disable_multi_reduction_to_contract_patterns }
}
@@ -68,7 +68,7 @@ func.func @contraction_batch_matmul(%A: memref<1584x1584x1584xf32>, %B: memref<1
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.batch_matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.batch_matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1 { disable_multi_reduction_to_contract_patterns }
}
@@ -108,7 +108,7 @@ func.func @vectorization_test(%A: memref<8x16xf32>, %B: memref<16x32xf32>,
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1 { disable_multi_reduction_to_contract_patterns, disable_transfer_permutation_map_lowering_patterns }
}
@@ -148,7 +148,7 @@ func.func @generic_output_transpose(%A: memref<8x16xf32>, %B: memref<16x32xf32>,
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1 { disable_multi_reduction_to_contract_patterns, disable_transfer_permutation_map_lowering_patterns }
}
@@ -175,7 +175,7 @@ func.func @generic_interchanged_transpose(%arg0: tensor<12x128x32xf32>) -> tenso
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1 { disable_multi_reduction_to_contract_patterns, disable_transfer_permutation_map_lowering_patterns }
}
@@ -215,7 +215,7 @@ func.func @vectorization_test_integer(%A: memref<8x16xi32>, %B: memref<16x32xi32
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1 { disable_multi_reduction_to_contract_patterns, disable_transfer_permutation_map_lowering_patterns }
}
@@ -235,7 +235,7 @@ func.func @vectorization_test_2(%A: memref<8x16xf32>, %B: memref<16x32xf32>,
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1 { disable_multi_reduction_to_contract_patterns }
}
@@ -259,7 +259,7 @@ func.func @test_vectorize_scalar_input(%A : memref<8x16xf32>, %arg0 : f32) {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1
}
@@ -283,7 +283,7 @@ func.func @test_do_not_vectorize_unsupported_element_types(%A : memref<8x16xcomp
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1
}
@@ -300,7 +300,7 @@ func.func @test_vectorize_fill(%A : memref<8x16xf32>, %arg0 : f32) {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.fill"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.fill"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1
}
@@ -318,7 +318,7 @@ func.func @test_vectorize_fill_scalar(%A : memref<f32>, %arg0 : f32) {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.fill"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.fill"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1
}
@@ -335,7 +335,7 @@ func.func @test_vectorize_copy(%A : memref<8x16xf32>, %B : memref<8x16xf32>) {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["memref.copy"]} in %arg1
+ %0 = transform.structured.match ops{["memref.copy"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1
}
@@ -355,7 +355,7 @@ func.func @test_vectorize_copy_scalar(%A : memref<f32>, %B : memref<f32>) {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["memref.copy"]} in %arg1
+ %0 = transform.structured.match ops{["memref.copy"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1
}
@@ -371,7 +371,7 @@ func.func @test_vectorize_copy_complex(%A : memref<8x16xcomplex<f32>>, %B : memr
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["memref.copy"]} in %arg1
+ %0 = transform.structured.match ops{["memref.copy"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1
}
@@ -399,7 +399,7 @@ func.func @test_vectorize_trailing_index(%arg0: memref<1x2x4x8xindex>) {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1
}
@@ -428,7 +428,7 @@ func.func @test_vectorize_inner_index(%arg0: memref<1x2x4x8xindex>) {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1
}
@@ -513,7 +513,7 @@ func.func @generic_vectorize(%arg0: memref<4x256xf32>,
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1 { disable_transfer_permutation_map_lowering_patterns }
}
@@ -604,7 +604,7 @@ func.func @generic_vectorize_tensor(%arg0: tensor<4x256xf32>,
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1 { disable_transfer_permutation_map_lowering_patterns }
}
@@ -648,7 +648,7 @@ func.func @generic_vectorize_broadcast_transpose(
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1 { disable_transfer_permutation_map_lowering_patterns }
}
@@ -691,7 +691,7 @@ func.func @vectorization_transpose(%A: memref<14x7xf32>, %B: memref<16x14xf32>,
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1 { disable_transfer_permutation_map_lowering_patterns }
}
@@ -723,7 +723,7 @@ func.func @matmul_tensors(
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1 { disable_multi_reduction_to_contract_patterns, disable_transfer_permutation_map_lowering_patterns }
}
@@ -752,7 +752,7 @@ func.func @pad_static(%arg0: tensor<2x?x2xf32>, %pad_value: f32) -> tensor<2x3x4
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.pad"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.pad"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1 { vectorize_padding }
}
@@ -781,7 +781,7 @@ func.func @pad_static_source(%arg0: tensor<2x5x2xf32>, %pad_value: f32) -> tenso
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.pad"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.pad"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1 { vectorize_padding }
}
@@ -818,7 +818,7 @@ func.func @pad_static_dynamic(%arg0: tensor<1x2x2x?xf32>, %low: index, %high: in
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.pad"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.pad"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1 { vectorize_padding }
}
@@ -838,7 +838,7 @@ func.func @pad_static_complex(%arg0: tensor<2x5x2xcomplex<f32>>, %pad_value: com
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.pad"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.pad"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1 { vectorize_padding }
}
@@ -868,7 +868,7 @@ func.func @pad_and_transfer_read(%arg0: tensor<5x6xf32>) -> vector<7x9xf32> {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.pad"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.pad"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1 { vectorize_padding }
}
@@ -901,7 +901,7 @@ func.func @pad_and_transfer_write_static(
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %3 = transform.structured.match ops{["tensor.pad"]} in %arg1
+ %3 = transform.structured.match ops{["tensor.pad"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
%5 = transform.structured.vectorize %4 { vectorize_padding }
}
@@ -938,7 +938,7 @@ func.func @pad_and_transfer_write_dynamic_static(
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %3 = transform.structured.match ops{["tensor.pad"]} in %arg1
+ %3 = transform.structured.match ops{["tensor.pad"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
%5 = transform.structured.vectorize %4 { vectorize_padding }
}
@@ -972,7 +972,7 @@ func.func @pad_and_insert_slice_source(
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %3 = transform.structured.match ops{["tensor.pad"]} in %arg1
+ %3 = transform.structured.match ops{["tensor.pad"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
%5 = transform.structured.vectorize %4 { vectorize_padding }
}
@@ -1000,7 +1000,7 @@ func.func @pad_and_insert_slice_dest(
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %3 = transform.structured.match ops{["tensor.pad"]} in %arg1
+ %3 = transform.structured.match ops{["tensor.pad"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
%5 = transform.structured.vectorize %4
}
@@ -1037,7 +1037,7 @@ func.func @pad_tensor_non_const_pad_value(%arg0: tensor<5x6xf32>) -> tensor<12x1
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %3 = transform.structured.match ops{["tensor.pad"]} in %arg1
+ %3 = transform.structured.match ops{["tensor.pad"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
%5 = transform.structured.vectorize %4 { vectorize_padding }
}
@@ -1072,7 +1072,7 @@ func.func @sum_exp(%input: tensor<4x16x8xf32>, %output: tensor<4x16xf32>)
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %3 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
%5 = transform.structured.vectorize %4
}
@@ -1117,7 +1117,7 @@ func.func @sum_exp_2(%input: tensor<3x2xf32>, %input_2: tensor<5x4xf32>, %output
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %3 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
%5 = transform.structured.vectorize %4 { disable_multi_reduction_to_contract_patterns, disable_transfer_permutation_map_lowering_patterns }
}
@@ -1147,7 +1147,7 @@ func.func @red_max_2d(%arg0: tensor<4x4xf32>) -> tensor<4xf32> {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %3 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
%5 = transform.structured.vectorize %4 { vectorize_padding }
}
@@ -1178,7 +1178,7 @@ func.func @red_min_2d(%arg0: tensor<4x4xf32>) -> tensor<4xf32> {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %3 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
%5 = transform.structured.vectorize %4
}
@@ -1208,7 +1208,7 @@ func.func @red_mul_2d(%arg0: tensor<4x4xf32>) -> tensor<4xf32> {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %3 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
%5 = transform.structured.vectorize %4
}
@@ -1238,7 +1238,7 @@ func.func @red_or_2d(%arg0: tensor<4x4xi1>) -> tensor<4xi1> {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %3 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
%5 = transform.structured.vectorize %4
}
@@ -1268,7 +1268,7 @@ func.func @red_and_2d(%arg0: tensor<4x4xi1>) -> tensor<4xi1> {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %3 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
%5 = transform.structured.vectorize %4
}
@@ -1298,7 +1298,7 @@ func.func @red_xor_2d(%arg0: tensor<4x4xi1>) -> tensor<4xi1> {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %3 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
%5 = transform.structured.vectorize %4
}
@@ -1332,7 +1332,7 @@ func.func @explicit_broadcast(%arg0: tensor<4x4xf32>, %arg1: tensor<4x1xf32>) ->
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %3 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
%5 = transform.structured.vectorize %4
}
@@ -1370,11 +1370,11 @@ func.func @fused_broadcast_red_2d(%arg0: tensor<4x4xf32>, %arg1: tensor<4x1xf32>
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.fill"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.fill"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1
- %3 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
%5 = transform.structured.vectorize %4
}
@@ -1417,7 +1417,7 @@ func.func @reduce_1d(%arg0: tensor<32xf32>) -> tensor<f32> {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1
}
@@ -1448,7 +1448,7 @@ func.func @not_projected_permutation(%arg0: tensor<8x8xf32>) -> tensor<6x6x3x3xf
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1
}
@@ -1487,7 +1487,7 @@ func.func @mixed_parallel_reduced_results(%arg0 : tensor<2x4x8xf32>,
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1 { disable_multi_reduction_to_contract_patterns, disable_transfer_permutation_map_lowering_patterns }
}
@@ -1523,7 +1523,7 @@ func.func @vectorize_1d_tensor_extract(%arg0: tensor<3xf32>, %arg1: tensor<4x3xi
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1
}
@@ -1560,7 +1560,7 @@ func.func @vectorize_nd_tensor_extract_constant_idx(%arg0: tensor<3x3xf32>, %arg
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1 { vectorize_nd_extract }
}
@@ -1596,7 +1596,7 @@ func.func @vectorize_nd_tensor_extract_idx_from_iteration_index(%arg0: tensor<3x
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1 { vectorize_nd_extract }
}
@@ -1644,7 +1644,7 @@ func.func @vectorize_nd_tensor_extract_index_from_tensor(%arg0: tensor<3x3xf32>,
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1 { vectorize_nd_extract }
}
@@ -1668,7 +1668,7 @@ func.func @vectorize_map(%arg0: memref<64xf32>,
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.map"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.map"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1
}
@@ -1687,7 +1687,7 @@ func.func @vectorize_transpose(%arg0: memref<16x32x64xf32>,
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.transpose"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.transpose"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1
}
@@ -1710,7 +1710,7 @@ func.func @vectorize_reduce(%arg0: memref<16x32x64xf32>,
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.reduce"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.reduce"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1
}
@@ -1745,7 +1745,7 @@ func.func @vectorize_dynamic_identity(%arg0: tensor<?xf32>,
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
transform.structured.masked_vectorize %0 vector_sizes [4]
}
@@ -1779,7 +1779,7 @@ func.func @vectorize_dynamic_1d_broadcast(%arg0: tensor<?xf32>,
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
transform.structured.masked_vectorize %0 vector_sizes [4]
}
@@ -1817,7 +1817,7 @@ func.func @vectorize_dynamic_2d_transpose(%arg0: tensor<?x?xf32>,
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
transform.structured.masked_vectorize %0 vector_sizes [4, 8]
}
@@ -1854,7 +1854,7 @@ func.func @vectorize_dynamic_generic_2d_broadcast(%arg0: tensor<?x?xf32>,
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
transform.structured.masked_vectorize %0 vector_sizes [4, 8]
}
@@ -1876,7 +1876,7 @@ func.func @vectorize_dynamic_reduction(%arg0: tensor<?x?xf32>,
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
transform.structured.masked_vectorize %0 vector_sizes [4, 8]
}
@@ -1914,7 +1914,7 @@ func.func @vectorize_dynamic_transpose_reduction(%arg0: tensor<?x?x?xf32>,
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
transform.structured.masked_vectorize %0 vector_sizes [4, 8, 16]
}
@@ -1962,7 +1962,7 @@ func.func @not_vectorizable(%arg0: tensor<1x?xf32>, %arg1: index, %arg2: index,
}
transform.sequence failures(propagate) {
^bb0(%arg0: !pdl.operation):
- %0 = transform.structured.match ops{["func.func"]} in %arg0
+ %0 = transform.structured.match ops{["func.func"]} in %arg0 : (!pdl.operation) -> !pdl.operation
%1 = transform.structured.vectorize %0
}
@@ -1996,7 +1996,7 @@ func.func @wrong_reduction_detection(%input: tensor<120x64xf32>) -> tensor<120x6
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
%2 = transform.structured.vectorize %1
}
diff --git a/mlir/test/Dialect/MemRef/transform-ops.mlir b/mlir/test/Dialect/MemRef/transform-ops.mlir
index 8a1b62832900f..97166809f8c06 100644
--- a/mlir/test/Dialect/MemRef/transform-ops.mlir
+++ b/mlir/test/Dialect/MemRef/transform-ops.mlir
@@ -30,7 +30,7 @@ func.func @multi_buffer(%in: memref<16xf32>) {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["memref.alloc"]} in %arg1
+ %0 = transform.structured.match ops{["memref.alloc"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = transform.memref.multibuffer %0 {factor = 2 : i64}
// Verify that the returned handle is usable.
transform.test_print_remark_at_operand %1, "transformed" : !pdl.operation
diff --git a/mlir/test/Dialect/SCF/transform-op-coalesce.mlir b/mlir/test/Dialect/SCF/transform-op-coalesce.mlir
index 4c84f62f9802b..2e7bbc53b196b 100644
--- a/mlir/test/Dialect/SCF/transform-op-coalesce.mlir
+++ b/mlir/test/Dialect/SCF/transform-op-coalesce.mlir
@@ -24,7 +24,7 @@ func.func @coalesce_inner() {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["scf.for"]} attributes {coalesce} in %arg1
+ %0 = transform.structured.match ops{["scf.for"]} attributes {coalesce} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = transform.cast %0 : !pdl.operation to !transform.op<"scf.for">
%2 = transform.loop.coalesce %1: (!transform.op<"scf.for">) -> (!transform.op<"scf.for">)
}
@@ -50,7 +50,7 @@ func.func @coalesce_outer(%arg1: memref<64x64xf32, 1>, %arg2: memref<64x64xf32,
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["affine.for"]} attributes {coalesce} in %arg1
+ %0 = transform.structured.match ops{["affine.for"]} attributes {coalesce} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = transform.cast %0 : !pdl.operation to !transform.op<"affine.for">
%2 = transform.loop.coalesce %1 : (!transform.op<"affine.for">) -> (!transform.op<"affine.for">)
}
@@ -85,7 +85,7 @@ func.func @coalesce_and_unroll(%arg1: memref<64x64xf32, 1>, %arg2: memref<64x64x
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["scf.for"]} attributes {coalesce} in %arg1
+ %0 = transform.structured.match ops{["scf.for"]} attributes {coalesce} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = transform.cast %0 : !pdl.operation to !transform.op<"scf.for">
%2 = transform.loop.coalesce %1 : (!transform.op<"scf.for">) -> (!transform.op<"scf.for">)
transform.loop.unroll %2 {factor = 3} : !transform.op<"scf.for">
diff --git a/mlir/test/Dialect/SCF/transform-ops-invalid.mlir b/mlir/test/Dialect/SCF/transform-ops-invalid.mlir
index 57812dee7eed9..da040fff82733 100644
--- a/mlir/test/Dialect/SCF/transform-ops-invalid.mlir
+++ b/mlir/test/Dialect/SCF/transform-ops-invalid.mlir
@@ -12,7 +12,7 @@ func.func @test_loops_do_not_get_coalesced() {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["affine.for"]} attributes {coalesce} in %arg1
+ %0 = transform.structured.match ops{["affine.for"]} attributes {coalesce} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = transform.cast %0 : !pdl.operation to !transform.op<"affine.for">
// expected-error @below {{failed to coalesce}}
%2 = transform.loop.coalesce %1: (!transform.op<"affine.for">) -> (!transform.op<"affine.for">)
@@ -29,7 +29,7 @@ func.func @test_loops_do_not_get_unrolled() {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["arith.addi"]} in %arg1
+ %0 = transform.structured.match ops{["arith.addi"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = transform.loop.get_parent_for %0 { affine = true } : (!pdl.operation) -> !transform.op<"affine.for">
// expected-error @below {{failed to unroll}}
transform.loop.unroll %1 { factor = 8 } : !transform.op<"affine.for">
@@ -55,7 +55,7 @@ func.func @loop_outline_op_multi_region() {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["scf.while"]} in %arg1
+ %0 = transform.structured.match ops{["scf.while"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{failed to outline}}
transform.loop.outline %0 {func_name = "foo"} : (!pdl.operation) -> !pdl.operation
}
diff --git a/mlir/test/Dialect/SCF/transform-ops.mlir b/mlir/test/Dialect/SCF/transform-ops.mlir
index 0e4b3846369d3..d876d0f6be9c2 100644
--- a/mlir/test/Dialect/SCF/transform-ops.mlir
+++ b/mlir/test/Dialect/SCF/transform-ops.mlir
@@ -17,7 +17,7 @@ func.func @get_parent_for_op(%arg0: index, %arg1: index, %arg2: index) {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["arith.addi"]} in %arg1
+ %0 = transform.structured.match ops{["arith.addi"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// CHECK: = transform.loop.get_parent_for
%1 = transform.loop.get_parent_for %0 : (!pdl.operation) -> !transform.op<"scf.for">
%2 = transform.loop.get_parent_for %0 { num_loops = 2 } : (!pdl.operation) -> !transform.op<"scf.for">
@@ -37,7 +37,7 @@ func.func @get_parent_for_op_no_loop(%arg0: index, %arg1: index) {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["arith.addi"]} in %arg1
+ %0 = transform.structured.match ops{["arith.addi"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{could not find an 'scf.for' parent}}
%1 = transform.loop.get_parent_for %0 : (!pdl.operation) -> !transform.op<"scf.for">
}
@@ -76,7 +76,7 @@ func.func @loop_outline_op(%arg0: index, %arg1: index, %arg2: index) {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["arith.addi"]} in %arg1
+ %0 = transform.structured.match ops{["arith.addi"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = transform.loop.get_parent_for %0 : (!pdl.operation) -> !transform.op<"scf.for">
// CHECK: = transform.loop.outline %{{.*}}
transform.loop.outline %1 {func_name = "foo"} : (!transform.op<"scf.for">) -> !pdl.operation
@@ -105,7 +105,7 @@ func.func @loop_peel_op() {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["arith.addi"]} in %arg1
+ %0 = transform.structured.match ops{["arith.addi"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = transform.loop.get_parent_for %0 : (!pdl.operation) -> !transform.op<"scf.for">
transform.loop.peel %1 : (!transform.op<"scf.for">) -> !pdl.operation
}
@@ -138,7 +138,7 @@ func.func @loop_pipeline_op(%A: memref<?xf32>, %result: memref<?xf32>) {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["arith.addf"]} in %arg1
+ %0 = transform.structured.match ops{["arith.addf"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = transform.loop.get_parent_for %0 : (!pdl.operation) -> !transform.op<"scf.for">
%2 = transform.loop.pipeline %1 : (!transform.op<"scf.for">) -> !pdl.operation
// Verify that the returned handle is usable.
@@ -162,7 +162,7 @@ func.func @loop_unroll_op() {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["arith.addi"]} in %arg1
+ %0 = transform.structured.match ops{["arith.addi"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = transform.loop.get_parent_for %0 : (!pdl.operation) -> !transform.op<"scf.for">
transform.loop.unroll %1 { factor = 4 } : !transform.op<"scf.for">
}
@@ -186,7 +186,7 @@ func.func @get_parent_for_op(%arg0: index, %arg1: index, %arg2: index) {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["arith.addi"]} in %arg1
+ %0 = transform.structured.match ops{["arith.addi"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// CHECK: = transform.loop.get_parent_for
%1 = transform.loop.get_parent_for %0 { affine = true } : (!pdl.operation) -> !transform.op<"affine.for">
%2 = transform.loop.get_parent_for %0 { num_loops = 2, affine = true } : (!pdl.operation) -> !transform.op<"affine.for">
@@ -206,7 +206,7 @@ func.func @get_parent_for_op_no_loop(%arg0: index, %arg1: index) {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["arith.addi"]} in %arg1
+ %0 = transform.structured.match ops{["arith.addi"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{could not find an 'affine.for' parent}}
%1 = transform.loop.get_parent_for %0 { affine = true } : (!pdl.operation) -> !transform.op<"affine.for">
}
@@ -228,7 +228,7 @@ func.func @loop_unroll_op() {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["arith.addi"]} in %arg1
+ %0 = transform.structured.match ops{["arith.addi"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = transform.loop.get_parent_for %0 { affine = true } : (!pdl.operation) -> !transform.op<"affine.for">
transform.test_print_remark_at_operand %1, "affine for loop" : !transform.op<"affine.for">
transform.loop.unroll %1 { factor = 4, affine = true } : !transform.op<"affine.for">
@@ -253,7 +253,7 @@ func.func @test_mixed_loops() {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["arith.addi"]} in %arg1
+ %0 = transform.structured.match ops{["arith.addi"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1 = transform.loop.get_parent_for %0 { num_loops = 1, affine = true } : (!pdl.operation) -> !transform.op<"affine.for">
transform.test_print_remark_at_operand %1, "affine for loop" : !transform.op<"affine.for">
transform.loop.unroll %1 { factor = 4 } : !transform.op<"affine.for">
diff --git a/mlir/test/Dialect/Tensor/tiling.mlir b/mlir/test/Dialect/Tensor/tiling.mlir
index c874666737f8c..79276b6bd0dae 100644
--- a/mlir/test/Dialect/Tensor/tiling.mlir
+++ b/mlir/test/Dialect/Tensor/tiling.mlir
@@ -33,7 +33,7 @@ func.func @dynamic_pad_tensor_3_4(%input_tensor: tensor<?x?xf32>,
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.pad"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.pad"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:2 = transform.structured.tile_to_scf_for %0 [2, 3]
}
@@ -70,7 +70,7 @@ func.func @dynamic_pad_tensor_0_3(%input_tensor: tensor<?x?xf32>,
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.pad"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.pad"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loop = transform.structured.tile_to_scf_for %0 [0, 3]
}
@@ -104,7 +104,7 @@ func.func @static_pad_tensor_3_4(%input_tensor: tensor<7x9xf32>,
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.pad"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.pad"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:2 = transform.structured.tile_to_scf_for %0 [2, 3]
}
@@ -136,7 +136,7 @@ func.func @static_pad_tensor_0_3(%input_tensor: tensor<7x9xf32>,
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.pad"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.pad"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loop = transform.structured.tile_to_scf_for %0 [0, 3]
}
@@ -174,7 +174,7 @@ func.func @static_pad_tile_evenly_0_3(%input_tensor: tensor<7x9xf32>,
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.pad"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.pad"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loop = transform.structured.tile_to_scf_for %0 [0, 3]
}
@@ -214,7 +214,7 @@ func.func @NC_to_NCnc(%arg0: tensor<128x256xf32>, %arg1: tensor<4x8x32x32xf32>)
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.pack"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.pack"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:2 = transform.structured.tile_to_scf_for %0 [2, 4]
}
@@ -244,7 +244,7 @@ func.func @KC_to_CKkc(%arg0: tensor<128x256xf32>, %arg1: tensor<32x4x32x8xf32>)
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.pack"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.pack"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:2 = transform.structured.tile_to_scf_for %0 [2, 4]
}
@@ -279,7 +279,7 @@ func.func @pad_and_pack_static(%input: tensor<13x15xf32>, %output: tensor<2x8x8x
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.pack"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.pack"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:2 = transform.structured.tile_to_scf_for %0 [2, 4]
}
@@ -328,7 +328,7 @@ func.func @pad_and_pack_partially_dynamic(%input: tensor<?x?xf32>, %output: tens
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.pack"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.pack"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:2 = transform.structured.tile_to_scf_for %0 [2, 4]
}
@@ -382,7 +382,7 @@ func.func @pad_and_pack_fully_dynamic(%source: tensor<?x?xf32>, %dest: tensor<?x
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.pack"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.pack"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:2 = transform.structured.tile_to_scf_for %0 [2, 4]
}
@@ -431,7 +431,7 @@ func.func @NCnc_to_NC(%source: tensor<8x8x32x16xf32>, %dest: tensor<256x128xf32>
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.unpack"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.unpack"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:2 = transform.structured.tile_to_scf_for %0 [2, 4]
}
@@ -479,7 +479,7 @@ func.func @CKkc_to_KC(%source: tensor<32x4x32x8xf32>, %dest: tensor<128x256xf32>
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.unpack"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.unpack"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:2 = transform.structured.tile_to_scf_for %0 [2, 4]
}
@@ -515,7 +515,7 @@ func.func @perfect_CKkc_to_KC(%source: tensor<32x4x2x4xf32>, %dest: tensor<8x128
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.unpack"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.unpack"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:2 = transform.structured.tile_to_scf_for %0 [2, 4]
}
@@ -557,7 +557,7 @@ func.func @dynamic_perfect_CKkc_to_KC(%source: tensor<?x?x2x2xf32>, %dest: tenso
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.unpack"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.unpack"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:2 = transform.structured.tile_to_scf_for %0 [2, 4]
}
@@ -592,7 +592,7 @@ func.func @perfect_NKPQk_to_NPQK(%source: tensor<1x4x6x6x2xf32>, %dest: tensor<1
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.unpack"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.unpack"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:4 = transform.structured.tile_to_scf_for %0 [1, 1, 1, 4]
}
@@ -631,6 +631,6 @@ func.func @perfect_NPQK_to_NKPQk(%source: tensor<1x6x6x8xf32>, %dest: tensor<1x4
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["tensor.pack"]} in %arg1
+ %0 = transform.structured.match ops{["tensor.pack"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:4 = transform.structured.tile_to_scf_for %0 [1, 1, 1, 1]
}
diff --git a/mlir/test/Dialect/Transform/test-interpreter.mlir b/mlir/test/Dialect/Transform/test-interpreter.mlir
index 1e99f9dfe8381..b0b16abc0399c 100644
--- a/mlir/test/Dialect/Transform/test-interpreter.mlir
+++ b/mlir/test/Dialect/Transform/test-interpreter.mlir
@@ -750,7 +750,7 @@ func.func @get_parent_for_op_no_loop(%arg0: index, %arg1: index) {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %addi = transform.structured.match ops{["arith.addi"]} in %arg1
+ %addi = transform.structured.match ops{["arith.addi"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%muli = get_producer_of_operand %addi[0] : (!pdl.operation) -> !pdl.operation
transform.test_print_remark_at_operand %muli, "found muli" : !pdl.operation
}
@@ -765,7 +765,7 @@ func.func @get_parent_for_op_no_loop(%arg0: index, %arg1: index) {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %muli = transform.structured.match ops{["arith.muli"]} in %arg1
+ %muli = transform.structured.match ops{["arith.muli"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{could not find a producer for operand number: 0 of}}
%bbarg = get_producer_of_operand %muli[0] : (!pdl.operation) -> !pdl.operation
@@ -782,7 +782,7 @@ func.func @get_consumer(%arg0: index, %arg1: index) {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %muli = transform.structured.match ops{["arith.muli"]} in %arg1
+ %muli = transform.structured.match ops{["arith.muli"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%addi = get_consumers_of_result %muli[0] : (!pdl.operation) -> !pdl.operation
transform.test_print_remark_at_operand %addi, "found addi" : !pdl.operation
}
@@ -797,7 +797,7 @@ func.func @get_consumer_fail_1(%arg0: index, %arg1: index) {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %muli = transform.structured.match ops{["arith.muli"]} in %arg1
+ %muli = transform.structured.match ops{["arith.muli"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{handle must be mapped to exactly one payload op}}
%bbarg = get_consumers_of_result %muli[0] : (!pdl.operation) -> !pdl.operation
@@ -812,7 +812,7 @@ func.func @get_consumer_fail_2(%arg0: index, %arg1: index) {
transform.sequence failures(propagate) {
^bb1(%arg1: !pdl.operation):
- %muli = transform.structured.match ops{["arith.muli"]} in %arg1
+ %muli = transform.structured.match ops{["arith.muli"]} in %arg1 : (!pdl.operation) -> !pdl.operation
// expected-error @below {{result number overflow}}
%bbarg = get_consumers_of_result %muli[1] : (!pdl.operation) -> !pdl.operation
@@ -828,11 +828,11 @@ func.func @split_handles(%a: index, %b: index, %c: index) {
transform.sequence failures(propagate) {
^bb1(%fun: !pdl.operation):
- %muli = transform.structured.match ops{["arith.muli"]} in %fun
+ %muli = transform.structured.match ops{["arith.muli"]} in %fun : (!pdl.operation) -> !pdl.operation
%h:2 = split_handles %muli in [2] : (!pdl.operation) -> (!pdl.operation, !pdl.operation)
// expected-remark @below {{1}}
transform.test_print_number_of_associated_payload_ir_ops %h#0
- %muli_2 = transform.structured.match ops{["arith.muli"]} in %fun
+ %muli_2 = transform.structured.match ops{["arith.muli"]} in %fun : (!pdl.operation) -> !pdl.operation
// expected-error @below {{expected to contain 3 operation handles but it only contains 2 handles}}
%h_2:3 = split_handles %muli_2 in [3] : (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation)
}
@@ -847,11 +847,11 @@ func.func @split_handles(%a: index, %b: index, %c: index) {
transform.sequence failures(suppress) {
^bb1(%fun: !pdl.operation):
- %muli = transform.structured.match ops{["arith.muli"]} in %fun
+ %muli = transform.structured.match ops{["arith.muli"]} in %fun : (!pdl.operation) -> !pdl.operation
%h:2 = split_handles %muli in [2] : (!pdl.operation) -> (!pdl.operation, !pdl.operation)
// expected-remark @below {{1}}
transform.test_print_number_of_associated_payload_ir_ops %h#0
- %muli_2 = transform.structured.match ops{["arith.muli"]} in %fun
+ %muli_2 = transform.structured.match ops{["arith.muli"]} in %fun : (!pdl.operation) -> !pdl.operation
// Silenceable failure and all handles are now empty.
%h_2:3 = split_handles %muli_2 in [3] : (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation)
// expected-remark @below {{0}}
@@ -976,7 +976,7 @@ func.func @split_handles(%a: index, %b: index, %c: index) {
transform.sequence -> !pdl.operation failures(propagate) {
^bb1(%fun: !pdl.operation):
- %muli = transform.structured.match ops{["arith.muli"]} in %fun
+ %muli = transform.structured.match ops{["arith.muli"]} in %fun : (!pdl.operation) -> !pdl.operation
// expected-error @below {{expected to contain 3 operation handles but it only contains 2 handles}}
%h_2:3 = split_handles %muli in [3] : (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation)
/// Test that yield does not crash in the presence of silenceable error in
@@ -1016,7 +1016,7 @@ transform.sequence failures(propagate) {
transform.sequence failures(propagate) {
^bb0(%arg0: !pdl.operation):
- %0 = transform.structured.match ops{["func.func"]} in %arg0
+ %0 = transform.structured.match ops{["func.func"]} in %arg0 : (!pdl.operation) -> !pdl.operation
%1 = transform.test_produce_param_with_number_of_test_ops %0 : !pdl.operation
// expected-remark @below {{1 : i32, 3 : i32}}
transform.test_print_param %1 : !transform.test_dialect_param
diff --git a/mlir/test/Dialect/Vector/transform-vector.mlir b/mlir/test/Dialect/Vector/transform-vector.mlir
index 864fd8ffc3476..cf3738f2e9b5e 100644
--- a/mlir/test/Dialect/Vector/transform-vector.mlir
+++ b/mlir/test/Dialect/Vector/transform-vector.mlir
@@ -15,12 +15,12 @@ func.func @matmul_tensors(
transform.sequence failures(propagate) {
^bb1(%module_op: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %module_op
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %module_op : (!pdl.operation) -> !pdl.operation
%1, %loops:3 = transform.structured.tile %0 [8, 4, 2] : (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation, !pdl.operation)
%2 = get_closest_isolated_parent %1 : (!pdl.operation) -> !pdl.operation
transform.structured.vectorize %2
transform.bufferization.one_shot_bufferize %module_op
- %func = transform.structured.match ops{["func.func"]} in %module_op
+ %func = transform.structured.match ops{["func.func"]} in %module_op : (!pdl.operation) -> !pdl.operation
transform.vector.lower_vectors %func multireduction_lowering = "innerreduction"
}
diff --git a/mlir/test/Integration/Dialect/Linalg/CPU/test-conv-1d-call.mlir b/mlir/test/Integration/Dialect/Linalg/CPU/test-conv-1d-call.mlir
index 6e66dec6a9fec..1505ccb439fe9 100644
--- a/mlir/test/Integration/Dialect/Linalg/CPU/test-conv-1d-call.mlir
+++ b/mlir/test/Integration/Dialect/Linalg/CPU/test-conv-1d-call.mlir
@@ -26,7 +26,7 @@ func.func @conv_1d(%arg0: memref<?xf32>, %arg1: memref<?xf32>, %arg2: memref<?xf
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.conv_1d"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.conv_1d"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loop = transform.structured.tile %0 [4] : (!pdl.operation) -> (!pdl.operation, !pdl.operation)
}
diff --git a/mlir/test/Integration/Dialect/Linalg/CPU/test-conv-1d-nwc-wcf-call.mlir b/mlir/test/Integration/Dialect/Linalg/CPU/test-conv-1d-nwc-wcf-call.mlir
index cbb55895b1ac9..df760b9ebd23f 100644
--- a/mlir/test/Integration/Dialect/Linalg/CPU/test-conv-1d-nwc-wcf-call.mlir
+++ b/mlir/test/Integration/Dialect/Linalg/CPU/test-conv-1d-nwc-wcf-call.mlir
@@ -28,7 +28,7 @@ func.func @conv_1d_nwc_wcf(%arg0: memref<?x?x?xf32>, %arg1: memref<?x?x?xf32>, %
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.conv_1d_nwc_wcf"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.conv_1d_nwc_wcf"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:2 = transform.structured.tile %0 [2, 4] : (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation)
}
diff --git a/mlir/test/Integration/Dialect/Linalg/CPU/test-conv-2d-call.mlir b/mlir/test/Integration/Dialect/Linalg/CPU/test-conv-2d-call.mlir
index ca5d0c9213d39..b39ed2cc5ebb3 100644
--- a/mlir/test/Integration/Dialect/Linalg/CPU/test-conv-2d-call.mlir
+++ b/mlir/test/Integration/Dialect/Linalg/CPU/test-conv-2d-call.mlir
@@ -26,7 +26,7 @@ func.func @conv_2d(%arg0: memref<?x?xf32>, %arg1: memref<?x?xf32>, %arg2: memref
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.conv_2d"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.conv_2d"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:2 = transform.structured.tile %0 [2, 2] : (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation)
}
diff --git a/mlir/test/Integration/Dialect/Linalg/CPU/test-conv-2d-nhwc-hwcf-call.mlir b/mlir/test/Integration/Dialect/Linalg/CPU/test-conv-2d-nhwc-hwcf-call.mlir
index 410de4bbd5758..365dd1713e140 100644
--- a/mlir/test/Integration/Dialect/Linalg/CPU/test-conv-2d-nhwc-hwcf-call.mlir
+++ b/mlir/test/Integration/Dialect/Linalg/CPU/test-conv-2d-nhwc-hwcf-call.mlir
@@ -28,7 +28,7 @@ func.func @conv_2d_nhwc_hwcf(%arg0: memref<?x?x?x?xf32>, %arg1: memref<?x?x?x?xf
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.conv_2d_nhwc_hwcf"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.conv_2d_nhwc_hwcf"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:4 = transform.structured.tile %0 [2, 3, 3, 2] : (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation, !pdl.operation, !pdl.operation)
}
diff --git a/mlir/test/Integration/Dialect/Linalg/CPU/test-conv-3d-call.mlir b/mlir/test/Integration/Dialect/Linalg/CPU/test-conv-3d-call.mlir
index a5a89a5d76067..45a7bdc1bd861 100644
--- a/mlir/test/Integration/Dialect/Linalg/CPU/test-conv-3d-call.mlir
+++ b/mlir/test/Integration/Dialect/Linalg/CPU/test-conv-3d-call.mlir
@@ -26,7 +26,7 @@ func.func @conv_3d(%arg0: memref<?x?x?xf32>, %arg1: memref<?x?x?xf32>, %arg2: me
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.conv_3d"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.conv_3d"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:3 = transform.structured.tile %0 [2, 2, 2] : (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation, !pdl.operation)
}
diff --git a/mlir/test/Integration/Dialect/Linalg/CPU/test-conv-3d-ndhwc-dhwcf-call.mlir b/mlir/test/Integration/Dialect/Linalg/CPU/test-conv-3d-ndhwc-dhwcf-call.mlir
index 4a3e35659d4ae..32383c592f1bb 100644
--- a/mlir/test/Integration/Dialect/Linalg/CPU/test-conv-3d-ndhwc-dhwcf-call.mlir
+++ b/mlir/test/Integration/Dialect/Linalg/CPU/test-conv-3d-ndhwc-dhwcf-call.mlir
@@ -28,7 +28,7 @@ func.func @conv_3d_ndhwc_dhwcf(%arg0: memref<?x?x?x?x?xf32>, %arg1: memref<?x?x?
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.conv_3d_ndhwc_dhwcf"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.conv_3d_ndhwc_dhwcf"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:3 = transform.structured.tile %0 [0, 5, 5, 5] : (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation, !pdl.operation)
}
diff --git a/mlir/test/Integration/Dialect/Linalg/CPU/test-tensor-matmul.mlir b/mlir/test/Integration/Dialect/Linalg/CPU/test-tensor-matmul.mlir
index 80fb41b96d11b..2c145a0bc7533 100644
--- a/mlir/test/Integration/Dialect/Linalg/CPU/test-tensor-matmul.mlir
+++ b/mlir/test/Integration/Dialect/Linalg/CPU/test-tensor-matmul.mlir
@@ -38,7 +38,7 @@ func.func @main() {
transform.sequence failures(propagate) {
^bb0(%arg1: !pdl.operation):
- %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
+ %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 : (!pdl.operation) -> !pdl.operation
%1, %loops:3 = transform.structured.tile %0 [1, 2, 3] : (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation, !pdl.operation)
}
More information about the Mlir-commits
mailing list