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 @@ -2,14 +2,11 @@ // Test One-Shot Bufferize. -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - sequence %arg0 : !pdl.operation failures(propagate) { - ^bb0(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["func.func"]} in %arg1 - transform.bufferization.one_shot_bufferize %0 - {target_is_module = false} - } +transform.sequence failures(propagate) { +^bb0(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["func.func"]} in %arg1 + transform.bufferization.one_shot_bufferize %0 + {target_is_module = false} } // CHECK-LABEL: func @test_function( @@ -34,14 +31,11 @@ // Test analysis of One-Shot Bufferize only. -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - sequence %arg0 : !pdl.operation failures(propagate) { - ^bb0(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["func.func"]} in %arg1 - transform.bufferization.one_shot_bufferize %0 - {target_is_module = false, test_analysis_only = true} - } +transform.sequence failures(propagate) { +^bb0(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["func.func"]} in %arg1 + transform.bufferization.one_shot_bufferize %0 + {target_is_module = false, test_analysis_only = true} } // CHECK-LABEL: func @test_function_analysis( @@ -60,14 +54,11 @@ // Test One-Shot Bufferize transform failure with an unknown op. This would be // allowed with `allow_unknown_ops`. -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - sequence %arg0 : !pdl.operation failures(propagate) { - ^bb0(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["func.func"]} in %arg1 - // expected-error @+1 {{bufferization failed}} - transform.bufferization.one_shot_bufferize %0 {target_is_module = false} - } +transform.sequence failures(propagate) { +^bb0(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["func.func"]} in %arg1 + // expected-error @+1 {{bufferization failed}} + transform.bufferization.one_shot_bufferize %0 {target_is_module = false} } func.func @test_unknown_op_failure() -> (tensor) { @@ -80,13 +71,10 @@ // Test One-Shot Bufferize transform failure with a module op. -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - sequence %arg0 : !pdl.operation failures(propagate) { - ^bb0(%arg1: !pdl.operation): - // %arg1 is the module - transform.bufferization.one_shot_bufferize %arg1 - } +transform.sequence failures(propagate) { +^bb0(%arg1: !pdl.operation): + // %arg1 is the module + transform.bufferization.one_shot_bufferize %arg1 } module { 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 @@ -44,15 +44,12 @@ return %y : memref<2 x 32 x f32> } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 - // 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] } - } +transform.sequence failures(propagate) { +^bb1(%arg0: !pdl.operation): + %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 + // 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,14 +85,12 @@ return %y : memref<2 x 32 x f32> } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 - // 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] } - } + +transform.sequence failures(propagate) { +^bb1(%arg0: !pdl.operation): + %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 + // 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] } } // ----- @@ -117,14 +112,11 @@ return %y : memref<2 x 32 x f32> } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 - // expected-error @below {{unsupported dynamic blockdim size}} - transform.gpu.map_nested_foreach_to_threads %funcop { blockDim = [128, 4, 1] } - } +transform.sequence failures(propagate) { +^bb1(%arg0: !pdl.operation): + %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 + // expected-error @below {{unsupported dynamic blockdim size}} + transform.gpu.map_nested_foreach_to_threads %funcop { blockDim = [128, 4, 1] } } // ----- @@ -145,14 +137,11 @@ return %y : memref<2x32x32x32xf32> } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 - // expected-error @below {{scf.foreach_thread with rank > 3 does not lower to gpu.thread_id}} - transform.gpu.map_nested_foreach_to_threads %funcop { blockDim = [128, 4, 1] } - } +transform.sequence failures(propagate) { +^bb1(%arg0: !pdl.operation): + %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 + // expected-error @below {{scf.foreach_thread with rank > 3 does not lower to gpu.thread_id}} + transform.gpu.map_nested_foreach_to_threads %funcop { blockDim = [128, 4, 1] } } // ----- @@ -168,16 +157,13 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %matmul = transform.structured.match ops{["linalg.matmul"]} in %arg0 - %foreach, %tiled = transform.structured.tile_to_foreach_thread_op %matmul num_threads [10, 20, 30] - %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 - // 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] } - } +transform.sequence failures(propagate) { +^bb1(%arg0: !pdl.operation): + %matmul = transform.structured.match ops{["linalg.matmul"]} in %arg0 + %foreach, %tiled = transform.structured.tile_to_foreach_thread_op %matmul num_threads [10, 20, 30] + %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 + // 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] } } // ----- 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 @@ -30,13 +30,10 @@ return %y : !type } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 - transform.gpu.map_foreach_to_blocks %funcop { blockDim = [12, 9, 1]} - } +transform.sequence failures(propagate) { +^bb1(%arg0: !pdl.operation): + %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 + transform.gpu.map_foreach_to_blocks %funcop { blockDim = [12, 9, 1]} } // ----- @@ -87,13 +84,10 @@ return %y : !type } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 - transform.gpu.map_nested_foreach_to_threads %funcop { blockDim = [12, 9, 1] } - } +transform.sequence failures(propagate) { +^bb1(%arg0: !pdl.operation): + %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 + transform.gpu.map_nested_foreach_to_threads %funcop { blockDim = [12, 9, 1] } } // ----- @@ -129,14 +123,11 @@ return %y : !type4d } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %funcop = transform.structured.match ops{["func.func"]} in %arg0 - %gpuLaunch = transform.gpu.map_foreach_to_blocks %funcop { generate_gpu_launch } - transform.gpu.map_nested_foreach_to_threads %gpuLaunch { blockDim = [32, 4, 1] } - } +transform.sequence failures(propagate) { +^bb1(%arg0: !pdl.operation): + %funcop = transform.structured.match ops{["func.func"]} in %arg0 + %gpuLaunch = transform.gpu.map_foreach_to_blocks %funcop { generate_gpu_launch } + transform.gpu.map_nested_foreach_to_threads %gpuLaunch { blockDim = [32, 4, 1] } } // ----- @@ -166,11 +157,8 @@ return %y : !type } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 - transform.gpu.map_nested_foreach_to_threads %funcop { blockDim = [12, 9, 1], syncAfterDistribute = false } - } +transform.sequence failures(propagate) { +^bb1(%arg0: !pdl.operation): + %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 + transform.gpu.map_nested_foreach_to_threads %funcop { blockDim = [12, 9, 1], syncAfterDistribute = false } } 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 @@ -1,22 +1,19 @@ // RUN: mlir-opt --test-transform-dialect-interpreter --canonicalize %s | FileCheck %s -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - // This implements a 2D multisize tiling with target sizes [3, 10]. - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %1:3 = transform.structured.multitile_sizes %0 { dimension = 0, target_size = 3} - %t:3 = transform.structured.multitile_sizes %0 { dimension = 1, target_size = 10} - %2:2 = transform.structured.split %0 after %1#2 { dimension = 0 } - %3:2 = transform.structured.tile %2#0 [%1#0] - %4:2 = transform.structured.tile %2#1 [%1#1] - %5 = merge_handles %3#0, %4#0 : !pdl.operation - %tt:3 = replicate num(%5) %t#0, %t#1, %t#2 : !pdl.operation, !pdl.operation, !pdl.operation, !pdl.operation - %6:2 = transform.structured.split %5 after %tt#2 { dimension = 1 } - transform.structured.tile %6#0 [0, %tt#0] - transform.structured.tile %6#1 [0, %tt#1] - } +// 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 + %1:3 = transform.structured.multitile_sizes %0 { dimension = 0, target_size = 3} + %t:3 = transform.structured.multitile_sizes %0 { dimension = 1, target_size = 10} + %2:2 = transform.structured.split %0 after %1#2 { dimension = 0 } + %3:2 = transform.structured.tile %2#0 [%1#0] + %4:2 = transform.structured.tile %2#1 [%1#1] + %5 = merge_handles %3#0, %4#0 : !pdl.operation + %tt:3 = replicate num(%5) %t#0, %t#1, %t#2 : !pdl.operation, !pdl.operation, !pdl.operation, !pdl.operation + %6:2 = transform.structured.split %5 after %tt#2 { dimension = 1 } + transform.structured.tile %6#0 [0, %tt#0] + transform.structured.tile %6#1 [0, %tt#1] } func.func private @elem(%arg0: f32, %arg1: index, %arg2: index) -> f32 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 @@ -66,13 +66,10 @@ // CHECK-NOT: memref.dealloc %[[tmpB]] : memref<48xi8> // CHECK-NOT: memref.dealloc %[[tmpC]] : memref<24xi8> -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - sequence %arg0 : !pdl.operation failures(propagate) { - ^bb0(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 - %1 = transform.structured.promote %0 { use_alloca } - } +transform.sequence failures(propagate) { +^bb0(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 + %1 = transform.structured.promote %0 { use_alloca } } // ----- @@ -139,16 +136,12 @@ // CHECK: memref.dealloc %[[tmpB_f64]] : memref<96xi8> // CHECK: memref.dealloc %[[tmpC_f64]] : memref<48xi8> -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - sequence %arg0 : !pdl.operation failures(propagate) { - ^bb0(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 - %1 = transform.structured.promote %0 - } +transform.sequence failures(propagate) { +^bb0(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 + %1 = transform.structured.promote %0 } - // ----- #map6 = affine_map<(d0, d1, d2) -> (d0, d2)> @@ -189,11 +182,8 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - sequence %arg0 : !pdl.operation failures(propagate) { - ^bb0(%arg1: !pdl.operation): - %0 = transform.structured.match interface{LinalgOp} in %arg1 - %1 = transform.structured.promote %0 - } +transform.sequence failures(propagate) { +^bb0(%arg1: !pdl.operation): + %0 = transform.structured.match interface{LinalgOp} in %arg1 + %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 @@ -31,12 +31,9 @@ // CHECK: memref.dealloc %[[A0]] // CHECK: memref.dealloc %[[A1]] -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - sequence %arg0 : !pdl.operation failures(propagate) { - ^bb0(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 - %1, %loops:3 = transform.structured.tile %0 [16, 16, 16] - %2 = transform.structured.promote %1 { operands_to_promote = [0, 2], force_full_tiles = [false, false] } - } +transform.sequence failures(propagate) { +^bb0(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 + %1, %loops:3 = transform.structured.tile %0 [16, 16, 16] + %2 = transform.structured.promote %1 { operands_to_promote = [0, 2], force_full_tiles = [false, false] } } 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 @@ -32,13 +32,10 @@ return %0 : tensor } - transform.with_pdl_patterns { - ^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 - %1:2 = transform.structured.tile_to_foreach_thread_op %0 num_threads [10, 20] (mapped to dims [1, 0]) - } + transform.sequence failures(propagate) { + ^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 + %1:2 = transform.structured.tile_to_foreach_thread_op %0 num_threads [10, 20] (mapped to dims [1, 0]) } } @@ -76,13 +73,10 @@ return %0 : tensor<100x300xf32> } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 - %1:2 = transform.structured.tile_to_foreach_thread_op %0 num_threads [10, 21] - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 + %1:2 = transform.structured.tile_to_foreach_thread_op %0 num_threads [10, 21] } @@ -120,13 +114,10 @@ return %0 : tensor } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 - %1:2 = transform.structured.tile_to_foreach_thread_op %0 tile_sizes [10, 20] - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 + %1:2 = transform.structured.tile_to_foreach_thread_op %0 tile_sizes [10, 20] } // ----- @@ -161,13 +152,10 @@ return %0 : tensor<100x300xf32> } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 - %1:2 = transform.structured.tile_to_foreach_thread_op %0 tile_sizes [10, 21] - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 + %1:2 = transform.structured.tile_to_foreach_thread_op %0 tile_sizes [10, 21] } // ----- @@ -186,13 +174,10 @@ return %result : tensor<4xf32> } - transform.with_pdl_patterns { - ^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %1:2 = transform.structured.tile_to_foreach_thread_op %0 num_threads [2] (mapped to dims [0]) - } + transform.sequence failures(propagate) { + ^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %1:2 = transform.structured.tile_to_foreach_thread_op %0 num_threads [2] (mapped to dims [0]) } } // CHECK-DAG: #[[$map0:.+]] = affine_map<(d0) -> (d0 * 2)> @@ -240,14 +225,11 @@ return %0 : tensor } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 - %sz = transform.structured.match ops{["test.dummy"]} in %arg1 - %1:2 = transform.structured.tile_to_foreach_thread_op %0 tile_sizes [%sz, 20] - } +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 + %1:2 = transform.structured.tile_to_foreach_thread_op %0 tile_sizes [%sz, 20] } // ----- @@ -298,13 +280,10 @@ return %res1, %res2 : tensor<100xf32>, tensor<100xf32> } - transform.with_pdl_patterns { - ^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %foreach_thread, %tiled_generic = transform.structured.tile_to_foreach_thread_op %0 num_threads [7] - } + transform.sequence failures(propagate) { + ^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %foreach_thread, %tiled_generic = transform.structured.tile_to_foreach_thread_op %0 num_threads [7] } // ----- @@ -355,14 +334,9 @@ return %res2, %res3 : tensor<300x100xf32>, tensor<300xf32> } - transform.with_pdl_patterns { - ^bb0(%IN_MAT1: !pdl.operation): - transform.sequence %IN_MAT1 : !pdl.operation failures(propagate) { - ^bb1(%IN_MAT2: !pdl.operation): - %0 = transform.structured.match ops{["linalg.generic"]} in %IN_MAT2 - %foreach_thread, %tiled_generic = transform.structured.tile_to_foreach_thread_op %0 num_threads [4] - } + transform.sequence failures(propagate) { + ^bb1(%IN_MAT2: !pdl.operation): + %0 = transform.structured.match ops{["linalg.generic"]} in %IN_MAT2 + %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 @@ -56,11 +56,8 @@ return %0: tensor<1x1x56x96xf32> } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match interface{LinalgOp} in %arg1 - %1 = transform.structured.decompose %0 - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match interface{LinalgOp} in %arg1 + %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 @@ -41,16 +41,13 @@ func.func @dummy2() { return } func.func @dummy3() { return } - transform.with_pdl_patterns { - ^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation 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 - - // linalg.fill is tileable. The op is tiled and fused. - transform.structured.fuse_into_containing_op %0 into %1 - } + 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 + + // linalg.fill is tileable. The op is tiled and fused. + transform.structured.fuse_into_containing_op %0 into %1 } } @@ -87,16 +84,13 @@ func.return %2 : tensor<64xf32> } - transform.with_pdl_patterns { - ^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation 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 + 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 - // tensor.empty is not tileable. The op is cloned and fused. - transform.structured.fuse_into_containing_op %0 into %1 - } + // tensor.empty is not tileable. The op is cloned 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 @@ -15,13 +15,10 @@ return %1 : tensor } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.elemwise_binary"]} in %arg1 - %1, %loops:2 = transform.structured.fuse %0 {tile_sizes = [32, 32], tile_interchange = [0, 1]} - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.elemwise_binary"]} in %arg1 + %1, %loops:2 = transform.structured.fuse %0 {tile_sizes = [32, 32], tile_interchange = [0, 1]} } // ----- @@ -45,15 +42,12 @@ return %1 : tensor } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.elemwise_binary"]} in %arg1 - %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 - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.elemwise_binary"]} in %arg1 + %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 } // ----- @@ -94,12 +88,9 @@ func.return %0 : tensor<12x25xf32> } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %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] - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %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] } 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 @@ -10,11 +10,8 @@ return %0 : tensor } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.elemwise_unary"]} in %arg1 - %1 = transform.structured.generalize %0 - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.elemwise_unary"]} in %arg1 + %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 @@ -18,13 +18,10 @@ return %0 : tensor } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 - transform.structured.interchange %0 { iterator_interchange = [1, 0]} - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 + transform.structured.interchange %0 { iterator_interchange = [1, 0]} } // ----- @@ -35,12 +32,9 @@ return %0 : tensor } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 - // expected-error @below {{transform applied to the wrong op kind}} - transform.structured.interchange %0 { iterator_interchange = [1, 0]} - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 + // 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 @@ -9,18 +9,15 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %match_name = transform.structured.match ops{["arith.constant"]} in %arg1 - transform.test_print_remark_at_operand %match_name, "matched op name" : !pdl.operation - transform.test_consume_operand %match_name +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %match_name = transform.structured.match ops{["arith.constant"]} in %arg1 + 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 - transform.test_print_remark_at_operand %match_attr, "matched attr name" : !pdl.operation - transform.test_consume_operand %match_attr - } + %match_attr = transform.structured.match ops{["arith.constant"]} attributes{my_attr} in %arg1 + transform.test_print_remark_at_operand %match_attr, "matched attr name" : !pdl.operation + transform.test_consume_operand %match_attr } // ----- @@ -32,15 +29,12 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %match_name = transform.structured.match - ops{["arith.constant"]} filter_result_type = f32 in %arg1 - transform.test_print_remark_at_operand %match_name, "matched op name" : !pdl.operation - transform.test_consume_operand %match_name - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %match_name = transform.structured.match + ops{["arith.constant"]} filter_result_type = f32 in %arg1 + transform.test_print_remark_at_operand %match_name, "matched op name" : !pdl.operation + transform.test_consume_operand %match_name } // ----- @@ -61,21 +55,18 @@ return %1 : tensor<128x12x32xf32> } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %match_attr = transform.structured.match - ops{["linalg.generic"]} - attributes{iterator_types = ["parallel", "parallel", "parallel"]} - in %arg1 - transform.test_print_remark_at_operand %match_attr, "matched complex attr" : !pdl.operation - transform.test_consume_operand %match_attr +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %match_attr = transform.structured.match + ops{["linalg.generic"]} + attributes{iterator_types = ["parallel", "parallel", "parallel"]} + in %arg1 + transform.test_print_remark_at_operand %match_attr, "matched complex attr" : !pdl.operation + transform.test_consume_operand %match_attr - %no_match = transform.structured.match - attributes{iterator_types = ["parallel", "parallel", "reduction"]} - in %arg1 - // expected-remark @below {{0}} - transform.test_print_number_of_associated_payload_ir_ops %no_match - } + %no_match = transform.structured.match + attributes{iterator_types = ["parallel", "parallel", "reduction"]} + in %arg1 +// 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 @@ -2,13 +2,10 @@ // CHECK-DAG: #[[$MAP13:.+]] = affine_map<() -> (13)> -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - sequence %arg0 : !pdl.operation failures(propagate) { - ^bb0(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 - transform.structured.multitile_sizes %0 { target_size = 3, dimension = 0 } - } +transform.sequence failures(propagate) { + ^bb0(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 + transform.structured.multitile_sizes %0 { target_size = 3, dimension = 0 } } // CHECK-LABEL: @multitile_sizes_static @@ -29,13 +26,10 @@ // ----- -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - sequence %arg0 : !pdl.operation failures(propagate) { - ^bb0(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 - transform.structured.multitile_sizes %0 { target_size = 3, divisor = 2, dimension = 0 } - } +transform.sequence failures(propagate) { + ^bb0(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 + transform.structured.multitile_sizes %0 { target_size = 3, divisor = 2, dimension = 0 } } // CHECK: #[[$MAP_A:.+]] = affine_map<()[s0] -> ([[A_IMPL:s0 floordiv 2]])> 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 @@ -31,13 +31,10 @@ func.return %5 : tensor<24x25xf32> } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 - %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]} - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 + %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]} } // ----- @@ -50,14 +47,11 @@ func.return %0 : tensor<24x25xf32> } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 - // 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]} - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 + // 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]} } // ----- @@ -70,14 +64,11 @@ func.return %0 : tensor<24x25xf32> } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 - // 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]} - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 + // 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]} } // ----- @@ -91,13 +82,10 @@ func.return %0 : tensor<24x25xf32> } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(suppress) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 - // 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]} - } +transform.sequence failures(suppress) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 + // 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-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 @@ -10,12 +10,9 @@ func.return %0 : tensor<24x25xf32> } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 - %1, %loops = transform.structured.tile %0 [10, 0, 0] - %2 = transform.structured.scalarize %1 - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 + %1, %loops = transform.structured.tile %0 [10, 0, 0] + %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 @@ -4,12 +4,12 @@ func.func @matmul_split(%A : tensor, %B: tensor<256x32xf32>, %C: tensor) -> tensor { // CHECK: bufferization.alloc_tensor({{.*}}) : tensor - // CHECK: linalg.generic + // CHECK: linalg.generic // CHECK-SAME: iterator_types = ["parallel", "parallel", "parallel", "reduction"] // CHECK-SAME: ins(%{{[a-zA-Z0-9]*}}, %{{[a-zA-Z0-9]*}}, %{{[a-zA-Z0-9]*}} : tensor, tensor<256x32xf32>, tensor<64x4xi1>) // CHECK-SAME: outs(%{{[a-zA-Z0-9]*}} : tensor) { - // CHECK: linalg.generic + // CHECK: linalg.generic // CHECK-SAME: iterator_types = ["parallel", "parallel", "reduction"] // CHECK-SAME: ins(%{{[a-zA-Z0-9]*}} : tensor) // CHECK-SAME: outs(%{{[a-zA-Z0-9]*}} : tensor) { @@ -18,12 +18,9 @@ return %0: tensor } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 - %1:4 = transform.structured.split_reduction %0 - { split_factor = 4, insert_split_dimension = 2, use_scaling_algorithm, use_alloc} - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 + %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 @@ -3,12 +3,12 @@ // CHECK-LABEL: func.func @matmul_split func.func @matmul_split(%A : tensor<16x256xf32>, %B: tensor<256x32xf32>, %C: tensor<16x32xf32>) -> tensor<16x32xf32> { - // CHECK: linalg.generic + // CHECK: linalg.generic // CHECK-SAME: iterator_types = ["parallel", "parallel", "parallel", "reduction"] // CHECK-SAME: ins(%{{[a-zA-Z0-9_]*}}, %{{[a-zA-Z0-9_]*}} : tensor<16x4x64xf32>, tensor<4x64x32xf32>) // CHECK-SAME: outs(%{{[a-zA-Z0-9_]*}} : tensor<16x32x4xf32>) { - // CHECK: linalg.generic + // CHECK: linalg.generic // CHECK-SAME: iterator_types = ["parallel", "parallel", "reduction"] // CHECK-SAME: ins(%{{[a-zA-Z0-9_]*}} : tensor<16x32x4xf32>) // CHECK-SAME: outs(%{{[a-zA-Z0-9_]*}} : tensor<16x32xf32>) { @@ -17,11 +17,8 @@ return %0: tensor<16x32xf32> } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 - %1:4 = transform.structured.split_reduction %0 { split_factor = 4, insert_split_dimension = 2} - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 + %1:4 = transform.structured.split_reduction %0 { split_factor = 4, insert_split_dimension = 2} } 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 @@ -1,12 +1,9 @@ // RUN: mlir-opt %s --test-transform-dialect-interpreter --split-input-file -verify-diagnostics | FileCheck %s -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %1:2 = transform.structured.split %0 after 42 { dimension = 0 } - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %1:2 = transform.structured.split %0 after 42 { dimension = 0 } } func.func private @elem(%arg0: f32, %arg1: index, %arg2: index) -> f32 @@ -74,14 +71,11 @@ // ----- -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %1 = transform.structured.match ops{["func.call"]} in %arg1 - transform.structured.split %0 after %1 { dimension = 0 } - } +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 + transform.structured.split %0 after %1 { dimension = 0 } } func.func private @get_size() -> index @@ -125,14 +119,11 @@ // ----- -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %1:2 = transform.structured.split %0 after 4 { dimension = 0} - %2:2 = transform.structured.split %1#1 after 16 { dimension = 1 } - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %1:2 = transform.structured.split %0 after 4 { dimension = 0} + %2:2 = transform.structured.split %1#1 after 16 { dimension = 1 } } func.func private @elem(%arg0: f32, %arg1: index, %arg2: index) -> f32 @@ -193,15 +184,12 @@ // ----- -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %1 = transform.structured.match ops{["func.call"]} in %arg1 - // 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 } - } +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 + // 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 } } func.func private @get_size() -> i64 @@ -222,15 +210,12 @@ // ----- -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %1 = transform.structured.match ops{["func.call"]} in %arg1 - // 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 } - } +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 + // 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 } } func.func private @get_size() -> i64 @@ -249,21 +234,11 @@ // ----- -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - pdl.pattern @func_return : benefit(1) { - %0 = pdl.operands - %1 = pdl.types - %2 = pdl.operation "func.return"(%0 : !pdl.range) -> (%1 : !pdl.range) - pdl.rewrite %2 with "transform.dialect" - } - - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["func.return"]} in %arg1 - // expected-error @below {{only applies to structured ops}} - transform.structured.split %0 after 16 { dimension = 1 } - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["func.return"]} in %arg1 + // expected-error @below {{only applies to structured ops}} + transform.structured.split %0 after 16 { dimension = 1 } } func.func @noop(%arg0: tensor<100xf32>, %arg1: tensor<100xf32>) -> tensor<100xf32> { @@ -273,21 +248,11 @@ // ----- -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - pdl.pattern @linalg_generic : benefit(1) { - %0 = pdl.operands - %1 = pdl.types - %2 = pdl.operation "linalg.generic"(%0 : !pdl.range) -> (%1 : !pdl.range) - pdl.rewrite %2 with "transform.dialect" - } - - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 - // expected-error @below {{dimension 1 does not exist in target op}} - transform.structured.split %0 after 16 { dimension = 1 } - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 + // expected-error @below {{dimension 1 does not exist in target op}} + transform.structured.split %0 after 16 { dimension = 1 } } func.func @one_d_static(%arg0: tensor<100xf32>, %arg1: tensor<100xf32>) -> tensor<100xf32> { 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 @@ -1,12 +1,9 @@ // RUN: mlir-opt --test-transform-dialect-interpreter --split-input-file %s | FileCheck %s -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - sequence %arg0 : !pdl.operation failures(propagate) { - ^bb0(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 - %1, %loops:3 = transform.structured.tile %0 [4, 4, 4] - } +transform.sequence failures(propagate) { +^bb0(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 + %1, %loops:3 = transform.structured.tile %0 [4, 4, 4] } // CHECK-LABEL: func @tile_linalg_matmul( @@ -39,14 +36,11 @@ // ----- -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - sequence %arg0 : !pdl.operation failures(propagate) { - ^bb0(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 - %1 = transform.structured.match ops{["func.call"]} in %arg1 - %2, %loops:3 = transform.structured.tile %0 [%1, %1, 4] - } +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 + %2, %loops:3 = transform.structured.tile %0 [%1, %1, 4] } func.func private @get_dynamic_tile_size() -> index 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 @@ -16,22 +16,11 @@ func.return %0 : tensor<24x25xf32> } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - pdl.pattern @pdl_target : benefit(1) { - %args = operands - %results = types - %0 = pdl.operation "linalg.matmul"(%args : !pdl.range) -> (%results : !pdl.range) - // TODO: we don't want this, but it is the required terminator for pdl.pattern - rewrite %0 with "transform.dialect" - } - - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 - %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation - %2 = transform.structured.vectorize %1 - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 + %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation + %2 = transform.structured.vectorize %1 } // ----- @@ -73,14 +62,11 @@ return %9 : tensor<24x25xf32> } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 - %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation - %2 = transform.structured.vectorize %1 - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 + %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation + %2 = transform.structured.vectorize %1 } // ----- @@ -124,14 +110,11 @@ return %9 : tensor<24x25xf32> } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 - %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation - %2 = transform.structured.vectorize %1 {vectorize_padding} - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 + %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation + %2 = transform.structured.vectorize %1 {vectorize_padding} } // ----- @@ -144,12 +127,9 @@ func.return %0 : tensor<24x25xf32> } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 - // expected-error @below {{op requires isolated-from-above targets}} - %2 = transform.structured.vectorize %0 - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 + // 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 @@ -142,13 +142,10 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 - transform.structured.interchange %0 { iterator_interchange = [1, 2, 0]} - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 + transform.structured.interchange %0 { iterator_interchange = [1, 2, 0]} } // CHECK-LABEL: func @permute_generic 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 @@ -58,13 +58,10 @@ // CHECK-SAME: ins(%[[v0]], %[[v1]] : memref, memref) // CHECK-SAME: outs(%[[v2]] : memref) -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - sequence %arg0 : !pdl.operation failures(propagate) { - ^bb0(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 - %1 = transform.structured.promote %0 { operands_to_promote = [0, 1, 2], use_full_tiles_by_default } - } +transform.sequence failures(propagate) { +^bb0(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 + %1 = transform.structured.promote %0 { operands_to_promote = [0, 1, 2], use_full_tiles_by_default } } // ----- 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 @@ -40,20 +40,17 @@ return %7 : tensor } - transform.with_pdl_patterns { - ^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation 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 + 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 - // Tile the root. - %foreach_thread_op, %tiled_op = transform.structured.tile_to_foreach_thread_op %root num_threads [10, 20] + // Tile the root. + %foreach_thread_op, %tiled_op = transform.structured.tile_to_foreach_thread_op %root num_threads [10, 20] - // Fuse all producers. - transform.structured.fuse_into_containing_op %producers into %foreach_thread_op - } + // Fuse all producers. + transform.structured.fuse_into_containing_op %producers into %foreach_thread_op } } @@ -100,20 +97,17 @@ return %7 : tensor } - transform.with_pdl_patterns { - ^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation 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 - %reversed_producers = transform.test_reverse_payload_ops %producers + 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 + %reversed_producers = transform.test_reverse_payload_ops %producers - // Tile the root. - %foreach_thread_op, %tiled_op = transform.structured.tile_to_foreach_thread_op %root num_threads [10, 20] + // Tile the root. + %foreach_thread_op, %tiled_op = transform.structured.tile_to_foreach_thread_op %root num_threads [10, 20] - // Fuse all producers. - transform.structured.fuse_into_containing_op %reversed_producers into %foreach_thread_op - } + // Fuse all producers. + transform.structured.fuse_into_containing_op %reversed_producers into %foreach_thread_op } } 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,14 +12,11 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.dot"]} in %arg1 - %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation - %2 = transform.structured.vectorize %1 { disable_multi_reduction_to_contract_patterns } - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.dot"]} in %arg1 + %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation + %2 = transform.structured.vectorize %1 { disable_multi_reduction_to_contract_patterns } } // ----- @@ -34,14 +31,11 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.matvec"]} in %arg1 - %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation - %2 = transform.structured.vectorize %1 { disable_multi_reduction_to_contract_patterns } - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.matvec"]} in %arg1 + %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation + %2 = transform.structured.vectorize %1 { disable_multi_reduction_to_contract_patterns } } // ----- @@ -55,14 +49,11 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 - %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation - %2 = transform.structured.vectorize %1 { disable_multi_reduction_to_contract_patterns } - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 + %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation + %2 = transform.structured.vectorize %1 { disable_multi_reduction_to_contract_patterns } } // ----- @@ -77,14 +68,11 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.batch_matmul"]} in %arg1 - %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation - %2 = transform.structured.vectorize %1 { disable_multi_reduction_to_contract_patterns } - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.batch_matmul"]} in %arg1 + %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation + %2 = transform.structured.vectorize %1 { disable_multi_reduction_to_contract_patterns } } // ----- @@ -120,14 +108,11 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %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 } - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %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 } } // ----- @@ -163,14 +148,11 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %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 } - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %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 } } // ----- @@ -193,14 +175,11 @@ return %1 : tensor<128x12x32xf32> } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %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 } - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %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 } } // ----- @@ -236,14 +215,11 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %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 } - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %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 } } // ----- @@ -259,14 +235,11 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 - %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation - %2 = transform.structured.vectorize %1 { disable_multi_reduction_to_contract_patterns } - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 + %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation + %2 = transform.structured.vectorize %1 { disable_multi_reduction_to_contract_patterns } } // ----- @@ -286,14 +259,11 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation - %2 = transform.structured.vectorize %1 - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation + %2 = transform.structured.vectorize %1 } // ----- @@ -313,14 +283,11 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation - %2 = transform.structured.vectorize %1 - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation + %2 = transform.structured.vectorize %1 } // ----- @@ -333,14 +300,11 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.fill"]} in %arg1 - %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation - %2 = transform.structured.vectorize %1 - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.fill"]} in %arg1 + %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation + %2 = transform.structured.vectorize %1 } // ----- @@ -354,14 +318,11 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.fill"]} in %arg1 - %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation - %2 = transform.structured.vectorize %1 - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.fill"]} in %arg1 + %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation + %2 = transform.structured.vectorize %1 } // ----- @@ -374,14 +335,11 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["memref.copy"]} in %arg1 - %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation - %2 = transform.structured.vectorize %1 - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["memref.copy"]} in %arg1 + %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation + %2 = transform.structured.vectorize %1 } // ----- @@ -397,14 +355,11 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["memref.copy"]} in %arg1 - %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation - %2 = transform.structured.vectorize %1 - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["memref.copy"]} in %arg1 + %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation + %2 = transform.structured.vectorize %1 } // ----- @@ -427,14 +382,11 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation - %2 = transform.structured.vectorize %1 - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation + %2 = transform.structured.vectorize %1 } // ----- @@ -459,14 +411,11 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation - %2 = transform.structured.vectorize %1 - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation + %2 = transform.structured.vectorize %1 } // ----- @@ -547,14 +496,11 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation - %2 = transform.structured.vectorize %1 { disable_transfer_permutation_map_lowering_patterns } - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation + %2 = transform.structured.vectorize %1 { disable_transfer_permutation_map_lowering_patterns } } // ----- @@ -641,14 +587,11 @@ tensor<4x256xf32>, tensor<4x256xf32>, tensor<4x256xf32> } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation - %2 = transform.structured.vectorize %1 { disable_transfer_permutation_map_lowering_patterns } - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation + %2 = transform.structured.vectorize %1 { disable_transfer_permutation_map_lowering_patterns } } // ----- @@ -688,14 +631,11 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation - %2 = transform.structured.vectorize %1 { disable_transfer_permutation_map_lowering_patterns } - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation + %2 = transform.structured.vectorize %1 { disable_transfer_permutation_map_lowering_patterns } } // ----- @@ -734,14 +674,11 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation - %2 = transform.structured.vectorize %1 { disable_transfer_permutation_map_lowering_patterns } - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation + %2 = transform.structured.vectorize %1 { disable_transfer_permutation_map_lowering_patterns } } // ----- @@ -769,14 +706,11 @@ return %0 : tensor<8x12xf32> } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 - %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 } - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1 + %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 } } // ----- @@ -801,14 +735,11 @@ } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["tensor.pad"]} in %arg1 - %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation - %2 = transform.structured.vectorize %1 { vectorize_padding } - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["tensor.pad"]} in %arg1 + %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation + %2 = transform.structured.vectorize %1 { vectorize_padding } } // ----- @@ -833,14 +764,11 @@ } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["tensor.pad"]} in %arg1 - %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation - %2 = transform.structured.vectorize %1 { vectorize_padding } - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["tensor.pad"]} in %arg1 + %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation + %2 = transform.structured.vectorize %1 { vectorize_padding } } @@ -873,14 +801,11 @@ } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["tensor.pad"]} in %arg1 - %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation - %2 = transform.structured.vectorize %1 { vectorize_padding } - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["tensor.pad"]} in %arg1 + %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation + %2 = transform.structured.vectorize %1 { vectorize_padding } } @@ -907,14 +832,11 @@ } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["tensor.pad"]} in %arg1 - %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation - %2 = transform.structured.vectorize %1 { vectorize_padding } - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["tensor.pad"]} in %arg1 + %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation + %2 = transform.structured.vectorize %1 { vectorize_padding } } // ----- @@ -943,14 +865,11 @@ return %3 : tensor<5x6xf32> } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %3 = transform.structured.match ops{["tensor.pad"]} in %arg1 - %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation - %5 = transform.structured.vectorize %4 { vectorize_padding } - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %3 = transform.structured.match ops{["tensor.pad"]} in %arg1 + %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation + %5 = transform.structured.vectorize %4 { vectorize_padding } } @@ -983,14 +902,11 @@ return %3 : tensor } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %3 = transform.structured.match ops{["tensor.pad"]} in %arg1 - %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation - %5 = transform.structured.vectorize %4 { vectorize_padding } - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %3 = transform.structured.match ops{["tensor.pad"]} in %arg1 + %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation + %5 = transform.structured.vectorize %4 { vectorize_padding } } @@ -1020,14 +936,11 @@ return %r : tensor<12x13xf32> } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %3 = transform.structured.match ops{["tensor.pad"]} in %arg1 - %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation - %5 = transform.structured.vectorize %4 { vectorize_padding } - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %3 = transform.structured.match ops{["tensor.pad"]} in %arg1 + %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation + %5 = transform.structured.vectorize %4 { vectorize_padding } } @@ -1051,14 +964,11 @@ return %r : tensor<1x12x13xf32> } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %3 = transform.structured.match ops{["tensor.pad"]} in %arg1 - %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation - %5 = transform.structured.vectorize %4 - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %3 = transform.structured.match ops{["tensor.pad"]} in %arg1 + %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation + %5 = transform.structured.vectorize %4 } // ----- @@ -1091,14 +1001,11 @@ } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %3 = transform.structured.match ops{["tensor.pad"]} in %arg1 - %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation - %5 = transform.structured.vectorize %4 { vectorize_padding } - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %3 = transform.structured.match ops{["tensor.pad"]} in %arg1 + %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation + %5 = transform.structured.vectorize %4 { vectorize_padding } } // ----- @@ -1129,14 +1036,11 @@ } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation - %5 = transform.structured.vectorize %4 - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation + %5 = transform.structured.vectorize %4 } // ----- @@ -1177,14 +1081,11 @@ } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %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 } - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %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 } } // ----- @@ -1210,14 +1111,11 @@ } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation - %5 = transform.structured.vectorize %4 { vectorize_padding } - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation + %5 = transform.structured.vectorize %4 { vectorize_padding } } // ----- @@ -1244,14 +1142,11 @@ } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation - %5 = transform.structured.vectorize %4 - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation + %5 = transform.structured.vectorize %4 } // ----- @@ -1277,14 +1172,11 @@ } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation - %5 = transform.structured.vectorize %4 - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation + %5 = transform.structured.vectorize %4 } // ----- @@ -1310,14 +1202,11 @@ } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation - %5 = transform.structured.vectorize %4 - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation + %5 = transform.structured.vectorize %4 } // ----- @@ -1343,14 +1232,11 @@ } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation - %5 = transform.structured.vectorize %4 - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation + %5 = transform.structured.vectorize %4 } // ----- @@ -1376,14 +1262,11 @@ } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation - %5 = transform.structured.vectorize %4 - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation + %5 = transform.structured.vectorize %4 } // ----- @@ -1413,14 +1296,11 @@ } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation - %5 = transform.structured.vectorize %4 - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation + %5 = transform.structured.vectorize %4 } // ----- @@ -1454,18 +1334,15 @@ } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.fill"]} in %arg1 - %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation - %2 = transform.structured.vectorize %1 +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.fill"]} in %arg1 + %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation + %2 = transform.structured.vectorize %1 - %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation - %5 = transform.structured.vectorize %4 - } + %3 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation + %5 = transform.structured.vectorize %4 } // ----- @@ -1504,14 +1381,11 @@ return %2 : tensor } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation - %2 = transform.structured.vectorize %1 - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation + %2 = transform.structured.vectorize %1 } @@ -1538,14 +1412,11 @@ return %result : tensor<6x6x3x3xf32> } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation - %2 = transform.structured.vectorize %1 - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation + %2 = transform.structured.vectorize %1 } // ----- @@ -1580,12 +1451,9 @@ // CHECK-DAG: vector.transfer_write %[[MUL]], %[[ARG2]] // CHECK-DAG: vector.transfer_write %[[ADD]], %[[ARG3]] -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - transform.sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 - %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 } - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 + %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 } } 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 @@ -15,19 +15,16 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["arith.addi"]} in %arg1 - // 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"> - %3 = transform.loop.get_parent_for %0 { num_loops = 3 } : (!pdl.operation) -> !transform.op<"scf.for"> - transform.test_print_remark_at_operand %1, "third loop" : !transform.op<"scf.for"> - transform.test_print_remark_at_operand %2, "second loop" : !transform.op<"scf.for"> - transform.test_print_remark_at_operand %3, "first loop" : !transform.op<"scf.for"> - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["arith.addi"]} in %arg1 + // 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"> + %3 = transform.loop.get_parent_for %0 { num_loops = 3 } : (!pdl.operation) -> !transform.op<"scf.for"> + transform.test_print_remark_at_operand %1, "third loop" : !transform.op<"scf.for"> + transform.test_print_remark_at_operand %2, "second loop" : !transform.op<"scf.for"> + transform.test_print_remark_at_operand %3, "first loop" : !transform.op<"scf.for"> } // ----- @@ -38,14 +35,11 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["arith.addi"]} in %arg1 - // expected-error @below {{could not find an 'scf.for' parent}} - %1 = transform.loop.get_parent_for %0 : (!pdl.operation) -> !transform.op<"scf.for"> - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["arith.addi"]} in %arg1 + // expected-error @below {{could not find an 'scf.for' parent}} + %1 = transform.loop.get_parent_for %0 : (!pdl.operation) -> !transform.op<"scf.for"> } // ----- @@ -80,15 +74,12 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["arith.addi"]} in %arg1 - %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 - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["arith.addi"]} in %arg1 + %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 } // ----- @@ -109,14 +100,11 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["scf.while"]} in %arg1 - // expected-error @below {{failed to outline}} - transform.loop.outline %0 {func_name = "foo"} : (!pdl.operation) -> !pdl.operation - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["scf.while"]} in %arg1 + // expected-error @below {{failed to outline}} + transform.loop.outline %0 {func_name = "foo"} : (!pdl.operation) -> !pdl.operation } // ----- @@ -140,14 +128,11 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["arith.addi"]} in %arg1 - %1 = transform.loop.get_parent_for %0 : (!pdl.operation) -> !transform.op<"scf.for"> - transform.loop.peel %1 : (!transform.op<"scf.for">) -> !pdl.operation - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["arith.addi"]} in %arg1 + %1 = transform.loop.get_parent_for %0 : (!pdl.operation) -> !transform.op<"scf.for"> + transform.loop.peel %1 : (!transform.op<"scf.for">) -> !pdl.operation } // ----- @@ -176,16 +161,13 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["arith.addf"]} in %arg1 - %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. - transform.test_print_remark_at_operand %2, "transformed" : !pdl.operation - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["arith.addf"]} in %arg1 + %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. + transform.test_print_remark_at_operand %2, "transformed" : !pdl.operation } // ----- @@ -203,13 +185,10 @@ return } -transform.with_pdl_patterns { -^bb0(%arg0: !pdl.operation): - sequence %arg0 : !pdl.operation failures(propagate) { - ^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["arith.addi"]} in %arg1 - %1 = transform.loop.get_parent_for %0 : (!pdl.operation) -> !transform.op<"scf.for"> - transform.loop.unroll %1 { factor = 4 } : !transform.op<"scf.for"> - } +transform.sequence failures(propagate) { +^bb1(%arg1: !pdl.operation): + %0 = transform.structured.match ops{["arith.addi"]} in %arg1 + %1 = transform.loop.get_parent_for %0 : (!pdl.operation) -> !transform.op<"scf.for"> + transform.loop.unroll %1 { factor = 4 } : !transform.op<"scf.for"> }