diff --git a/mlir/include/mlir/Dialect/Linalg/TransformOps/LinalgTransformOps.td b/mlir/include/mlir/Dialect/Linalg/TransformOps/LinalgTransformOps.td --- a/mlir/include/mlir/Dialect/Linalg/TransformOps/LinalgTransformOps.td +++ b/mlir/include/mlir/Dialect/Linalg/TransformOps/LinalgTransformOps.td @@ -260,13 +260,13 @@ it is a navigation op. }]; - let arguments = (ins PDL_Operation:$target, + let arguments = (ins TransformHandleTypeInterface:$target, OptionalAttr:$ops, OptionalAttr:$interface, OptionalAttr:$op_attrs, OptionalAttr:$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":$opNames)> @@ -278,6 +278,7 @@ (`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 --- 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 @@ 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 @@ 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) { ^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@+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) { ^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@+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 --- 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 @@ 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 @@ 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 @@ 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 --- a/mlir/test/Dialect/GPU/transform-gpu-failing.mlir +++ b/mlir/test/Dialect/GPU/transform-gpu-failing.mlir @@ -6,7 +6,7 @@ } 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 @@ } 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 @@ 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 @@ 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 @@ 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, #gpu.thread, #gpu.thread ] ) - %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 @@ } 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 @@ 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 @@ 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 @@ 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 @@ 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 is duplicated, cannot map different loops to the same processor}} transform.gpu.map_nested_foreach_to_threads %funcop { blockDim = [32, 32]} } @@ -296,7 +296,7 @@ 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, #gpu.thread, #gpu.thread}} 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 --- a/mlir/test/Dialect/GPU/transform-gpu.mlir +++ b/mlir/test/Dialect/GPU/transform-gpu.mlir @@ -32,7 +32,7 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 --- a/mlir/test/Dialect/LLVM/transform-e2e.mlir +++ b/mlir/test/Dialect/LLVM/transform-e2e.mlir @@ -14,12 +14,12 @@ 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 --- a/mlir/test/Dialect/Linalg/generalize-tensor-pack.mlir +++ b/mlir/test/Dialect/Linalg/generalize-tensor-pack.mlir @@ -93,7 +93,7 @@ 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 @@ 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 @@ // 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 --- a/mlir/test/Dialect/Linalg/generalize-tensor-unpack.mlir +++ b/mlir/test/Dialect/Linalg/generalize-tensor-unpack.mlir @@ -67,7 +67,7 @@ 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 @@ 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 @@ 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 --- 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 @@ 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 %t:3 = transform.structured.multitile_sizes %0 { dimension = 1, target_size = 10} : (!pdl.operation) -> !transform.param %2:2 = transform.structured.split %0 after %1#2 { dimension = 0 } : !pdl.operation, !transform.param diff --git a/mlir/test/Dialect/Linalg/promote.mlir b/mlir/test/Dialect/Linalg/promote.mlir --- a/mlir/test/Dialect/Linalg/promote.mlir +++ b/mlir/test/Dialect/Linalg/promote.mlir @@ -68,7 +68,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 = transform.structured.promote %0 { use_alloca } } @@ -138,7 +138,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 = transform.structured.promote %0 } @@ -184,6 +184,6 @@ 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 --- a/mlir/test/Dialect/Linalg/promotion_options.mlir +++ b/mlir/test/Dialect/Linalg/promotion_options.mlir @@ -36,7 +36,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 [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 --- a/mlir/test/Dialect/Linalg/tile-conv.mlir +++ b/mlir/test/Dialect/Linalg/tile-conv.mlir @@ -11,7 +11,7 @@ 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 --- a/mlir/test/Dialect/Linalg/tile-indexed.mlir +++ b/mlir/test/Dialect/Linalg/tile-indexed.mlir @@ -13,7 +13,7 @@ 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 @@ 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 --- a/mlir/test/Dialect/Linalg/tile-tensors.mlir +++ b/mlir/test/Dialect/Linalg/tile-tensors.mlir @@ -29,7 +29,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 [2, 3, 4] : (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation, !pdl.operation) } @@ -60,7 +60,7 @@ 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 @@ 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 --- a/mlir/test/Dialect/Linalg/tile-to-foreach-thread.mlir +++ b/mlir/test/Dialect/Linalg/tile-to-foreach-thread.mlir @@ -34,7 +34,7 @@ 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, #gpu.thread ] ) } } @@ -76,8 +76,8 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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]) } } @@ -269,8 +269,8 @@ 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) { ^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) { ^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 --- a/mlir/test/Dialect/Linalg/transform-op-decompose.mlir +++ b/mlir/test/Dialect/Linalg/transform-op-decompose.mlir @@ -184,6 +184,6 @@ 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 --- 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 --- a/mlir/test/Dialect/Linalg/transform-op-fuse.mlir +++ b/mlir/test/Dialect/Linalg/transform-op-fuse.mlir @@ -17,7 +17,7 @@ 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 @@ 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 @@ 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 @@ 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 --- a/mlir/test/Dialect/Linalg/transform-op-generalize.mlir +++ b/mlir/test/Dialect/Linalg/transform-op-generalize.mlir @@ -12,6 +12,6 @@ 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 --- a/mlir/test/Dialect/Linalg/transform-op-interchange.mlir +++ b/mlir/test/Dialect/Linalg/transform-op-interchange.mlir @@ -20,7 +20,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 transform.structured.interchange %0 iterator_interchange = [1, 0] } @@ -34,7 +34,7 @@ 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 --- a/mlir/test/Dialect/Linalg/transform-op-match.mlir +++ b/mlir/test/Dialect/Linalg/transform-op-match.mlir @@ -11,11 +11,11 @@ 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 @@ 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 @@ #linalg.iterator_type, #linalg.iterator_type, #linalg.iterator_type]} - 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 @@ #linalg.iterator_type, #linalg.iterator_type, #linalg.iterator_type]} - 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 --- 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 @@ 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 @@ -55,7 +55,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, divisor = 2, dimension = 0 } : (!pdl.operation) -> !pdl.operation } @@ -96,7 +96,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 // 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 diff --git a/mlir/test/Dialect/Linalg/transform-op-pack.mlir b/mlir/test/Dialect/Linalg/transform-op-pack.mlir --- a/mlir/test/Dialect/Linalg/transform-op-pack.mlir +++ b/mlir/test/Dialect/Linalg/transform-op-pack.mlir @@ -35,7 +35,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 transform.structured.pack %0 packed_sizes = [0, 4] : (!pdl.operation) -> (!transform.op<"linalg.generic">) } @@ -77,7 +77,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 = 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 --- a/mlir/test/Dialect/Linalg/transform-op-pad.mlir +++ b/mlir/test/Dialect/Linalg/transform-op-pad.mlir @@ -33,7 +33,7 @@ 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 @@ 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 @@ 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 @@ 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 --- a/mlir/test/Dialect/Linalg/transform-op-replace.mlir +++ b/mlir/test/Dialect/Linalg/transform-op-replace.mlir @@ -10,7 +10,7 @@ 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 @@ 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 @@ 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 --- a/mlir/test/Dialect/Linalg/transform-op-scalarize.mlir +++ b/mlir/test/Dialect/Linalg/transform-op-scalarize.mlir @@ -20,7 +20,7 @@ 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 --- 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 @@ 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 --- a/mlir/test/Dialect/Linalg/transform-op-split-reduction.mlir +++ b/mlir/test/Dialect/Linalg/transform-op-split-reduction.mlir @@ -33,7 +33,7 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 --- 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 @@ 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 @@ 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 @@ 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) { ^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 @@ 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 @@ 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 @@ 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 @@ 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 --- 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 @@ 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 @@ 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 // expected-error @below {{expected as many parameter values (0) as target ops (2)}} @@ -103,9 +103,9 @@ 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 --- a/mlir/test/Dialect/Linalg/transform-op-vectorize.mlir +++ b/mlir/test/Dialect/Linalg/transform-op-vectorize.mlir @@ -18,7 +18,7 @@ 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 @@ 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 @@ 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 @@ 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 --- a/mlir/test/Dialect/Linalg/transform-patterns.mlir +++ b/mlir/test/Dialect/Linalg/transform-patterns.mlir @@ -13,7 +13,7 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 --- a/mlir/test/Dialect/Linalg/transform-promotion.mlir +++ b/mlir/test/Dialect/Linalg/transform-promotion.mlir @@ -60,7 +60,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 = transform.structured.promote %0 { operands_to_promote = [0, 1, 2], use_full_tiles_by_default } } @@ -122,7 +122,7 @@ ^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 @@ ^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 @@ ^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 --- a/mlir/test/Dialect/Linalg/transform-tile-and-fuse.mlir +++ b/mlir/test/Dialect/Linalg/transform-tile-and-fuse.mlir @@ -43,8 +43,8 @@ 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 @@ 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 --- a/mlir/test/Dialect/Linalg/transform-tile-reduction.mlir +++ b/mlir/test/Dialect/Linalg/transform-tile-reduction.mlir @@ -16,7 +16,7 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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] } @@ -287,7 +287,7 @@ 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] @@ -324,7 +324,7 @@ 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] diff --git a/mlir/test/Dialect/Linalg/vectorization.mlir b/mlir/test/Dialect/Linalg/vectorization.mlir --- a/mlir/test/Dialect/Linalg/vectorization.mlir +++ b/mlir/test/Dialect/Linalg/vectorization.mlir @@ -12,7 +12,7 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ } 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 @@ 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 --- a/mlir/test/Dialect/MemRef/transform-ops.mlir +++ b/mlir/test/Dialect/MemRef/transform-ops.mlir @@ -30,7 +30,7 @@ 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 --- a/mlir/test/Dialect/SCF/transform-op-coalesce.mlir +++ b/mlir/test/Dialect/SCF/transform-op-coalesce.mlir @@ -24,7 +24,7 @@ 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 @@ 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 @@ 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 --- a/mlir/test/Dialect/SCF/transform-ops-invalid.mlir +++ b/mlir/test/Dialect/SCF/transform-ops-invalid.mlir @@ -12,7 +12,7 @@ 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 @@ 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 @@ 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 --- a/mlir/test/Dialect/SCF/transform-ops.mlir +++ b/mlir/test/Dialect/SCF/transform-ops.mlir @@ -17,7 +17,7 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 --- a/mlir/test/Dialect/Tensor/tiling.mlir +++ b/mlir/test/Dialect/Tensor/tiling.mlir @@ -33,7 +33,7 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 --- a/mlir/test/Dialect/Transform/test-interpreter.mlir +++ b/mlir/test/Dialect/Transform/test-interpreter.mlir @@ -750,7 +750,7 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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) { ^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 --- a/mlir/test/Dialect/Vector/transform-vector.mlir +++ b/mlir/test/Dialect/Vector/transform-vector.mlir @@ -15,12 +15,12 @@ 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 --- 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 @@ 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 --- 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 @@ 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 --- 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 @@ 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 --- 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 @@ 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 --- 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 @@ 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 --- 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 @@ 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 --- 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 @@ 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) }