diff options
Diffstat (limited to 'mlir')
26 files changed, 193 insertions, 212 deletions
diff --git a/mlir/include/mlir/Dialect/Affine/TransformOps/AffineTransformOps.h b/mlir/include/mlir/Dialect/Affine/TransformOps/AffineTransformOps.h index 4a27f3035e6b..fb6f96ec684c 100644 --- a/mlir/include/mlir/Dialect/Affine/TransformOps/AffineTransformOps.h +++ b/mlir/include/mlir/Dialect/Affine/TransformOps/AffineTransformOps.h @@ -9,7 +9,6 @@ #ifndef MLIR_DIALECT_AFFINE_TRANSFORMOPS_AFFINETRANSFORMOPS_H #define MLIR_DIALECT_AFFINE_TRANSFORMOPS_AFFINETRANSFORMOPS_H -#include "mlir/Dialect/PDL/IR/PDLTypes.h" #include "mlir/Dialect/Transform/IR/TransformInterfaces.h" #include "mlir/Dialect/Transform/IR/TransformTypes.h" #include "mlir/IR/OpImplementation.h" diff --git a/mlir/include/mlir/Dialect/Affine/TransformOps/AffineTransformOps.td b/mlir/include/mlir/Dialect/Affine/TransformOps/AffineTransformOps.td index 76f0d8e6439c..b74e4af6eedd 100644 --- a/mlir/include/mlir/Dialect/Affine/TransformOps/AffineTransformOps.td +++ b/mlir/include/mlir/Dialect/Affine/TransformOps/AffineTransformOps.td @@ -9,7 +9,6 @@ #ifndef AFFINE_TRANSFORM_OPS #define AFFINE_TRANSFORM_OPS -include "mlir/Dialect/PDL/IR/PDLTypes.td" include "mlir/Dialect/Transform/IR/TransformDialect.td" include "mlir/Dialect/Transform/IR/TransformInterfaces.td" include "mlir/Dialect/Transform/IR/TransformTypes.td" @@ -50,16 +49,17 @@ def SimplifyBoundedAffineOpsOp TODO: Allow mixed PDL_Operation/int64_t for lower_bounds and upper_bounds. }]; - let arguments = (ins PDL_Operation:$target, - Variadic<PDL_Operation>:$bounded_values, + let arguments = (ins TransformHandleTypeInterface:$target, + Variadic<TransformHandleTypeInterface>:$bounded_values, DenseI64ArrayAttr:$lower_bounds, DenseI64ArrayAttr:$upper_bounds); let results = (outs); let hasVerifier = 1; let assemblyFormat = [{ - $target `with` `[` $bounded_values `]` + $target `with` `[` ($bounded_values^ `:` type($bounded_values))? `]` `within` $lower_bounds `and` $upper_bounds attr-dict + `:` type($target) }]; } diff --git a/mlir/include/mlir/Dialect/Bufferization/TransformOps/BufferizationTransformOps.h b/mlir/include/mlir/Dialect/Bufferization/TransformOps/BufferizationTransformOps.h index 06204b6b879a..426e9d02efda 100644 --- a/mlir/include/mlir/Dialect/Bufferization/TransformOps/BufferizationTransformOps.h +++ b/mlir/include/mlir/Dialect/Bufferization/TransformOps/BufferizationTransformOps.h @@ -10,7 +10,6 @@ #define MLIR_DIALECT_BUFFERIZATION_TRANSFORMOPS_BUFFERIZATIONTRANSFORMOPS_H #include "mlir/Dialect/Bufferization/IR/BufferizableOpInterface.h" -#include "mlir/Dialect/PDL/IR/PDLTypes.h" #include "mlir/Dialect/Transform/IR/TransformInterfaces.h" #include "mlir/Dialect/Transform/IR/TransformTypes.h" #include "mlir/IR/OpImplementation.h" diff --git a/mlir/include/mlir/Dialect/Bufferization/TransformOps/BufferizationTransformOps.td b/mlir/include/mlir/Dialect/Bufferization/TransformOps/BufferizationTransformOps.td index 692f0f0b68d6..49b5ef6f0fee 100644 --- a/mlir/include/mlir/Dialect/Bufferization/TransformOps/BufferizationTransformOps.td +++ b/mlir/include/mlir/Dialect/Bufferization/TransformOps/BufferizationTransformOps.td @@ -13,7 +13,6 @@ include "mlir/Dialect/Bufferization/IR/BufferizationEnums.td" include "mlir/Dialect/Transform/IR/TransformDialect.td" include "mlir/Dialect/Transform/IR/TransformInterfaces.td" include "mlir/Dialect/Transform/IR/TransformTypes.td" -include "mlir/Dialect/PDL/IR/PDLTypes.td" include "mlir/Interfaces/SideEffectInterfaces.td" include "mlir/IR/OpBase.td" @@ -119,11 +118,11 @@ def EliminateEmptyTensorsOp not produce any handle. }]; - let arguments = (ins PDL_Operation:$target); + let arguments = (ins TransformHandleTypeInterface:$target); let results = (outs); - let assemblyFormat = "$target attr-dict"; + let assemblyFormat = "$target attr-dict `:` type($target)"; } //===----------------------------------------------------------------------===// diff --git a/mlir/include/mlir/Dialect/GPU/TransformOps/GPUTransformOps.h b/mlir/include/mlir/Dialect/GPU/TransformOps/GPUTransformOps.h index 57d74d856cba..3b10fcb77aaf 100644 --- a/mlir/include/mlir/Dialect/GPU/TransformOps/GPUTransformOps.h +++ b/mlir/include/mlir/Dialect/GPU/TransformOps/GPUTransformOps.h @@ -9,7 +9,6 @@ #ifndef MLIR_DIALECT_GPU_TRANSFORMOPS_GPUTRANSFORMOPS_H #define MLIR_DIALECT_GPU_TRANSFORMOPS_GPUTRANSFORMOPS_H -#include "mlir/Dialect/PDL/IR/PDLTypes.h" #include "mlir/Dialect/SCF/IR/SCF.h" #include "mlir/Dialect/Transform/IR/TransformInterfaces.h" #include "mlir/IR/OpImplementation.h" diff --git a/mlir/include/mlir/Dialect/GPU/TransformOps/GPUTransformOps.td b/mlir/include/mlir/Dialect/GPU/TransformOps/GPUTransformOps.td index c719fedc90e3..42162315aebd 100644 --- a/mlir/include/mlir/Dialect/GPU/TransformOps/GPUTransformOps.td +++ b/mlir/include/mlir/Dialect/GPU/TransformOps/GPUTransformOps.td @@ -11,7 +11,6 @@ include "mlir/Dialect/Transform/IR/TransformDialect.td" include "mlir/Dialect/Transform/IR/TransformInterfaces.td" -include "mlir/Dialect/PDL/IR/PDLTypes.td" include "mlir/Interfaces/SideEffectInterfaces.td" include "mlir/IR/OpBase.td" @@ -41,7 +40,7 @@ def MapNestedForallToThreads : to constrain the number of warps in each dimension. When present, this attribute must be specified in a way that is compatible with the block_dims attribute. If an `scf.forall` op is mapped to fewer warps, - predicaiton occurs. + predication occurs. Dynamic `scf.forall` trip counts are currently not supported. Dynamic block dim sizes are currently not supported. @@ -111,11 +110,11 @@ def MapNestedForallToThreads : ``` }]; - let arguments = (ins PDL_Operation:$target, + let arguments = (ins TransformHandleTypeInterface:$target, DefaultValuedAttr<DenseI64ArrayAttr, "{}">:$block_dims, DefaultValuedOptionalAttr<DenseI64ArrayAttr, "{}">:$warp_dims, DefaultValuedAttr<BoolAttr, "true">:$sync_after_distribute); - let results = (outs PDL_Operation:$result); + let results = (outs TransformHandleTypeInterface:$result); let assemblyFormat = [{ $target @@ -123,6 +122,7 @@ def MapNestedForallToThreads : (`warp_dims` `=` $warp_dims^)? (`sync_after_distribute` `=` $sync_after_distribute^)? attr-dict + `:` functional-type($target, $result) }]; let extraClassDeclaration = [{ ::mlir::DiagnosedSilenceableFailure applyToOne( @@ -175,16 +175,17 @@ def MapForallToBlocks : properties. }]; - let arguments = (ins PDL_Operation:$target, + let arguments = (ins TransformHandleTypeInterface:$target, DefaultValuedOptionalAttr<DenseI64ArrayAttr, "{}">:$grid_dims, UnitAttr:$generate_gpu_launch); - let results = (outs PDL_Operation:$result); + let results = (outs TransformHandleTypeInterface:$result); let assemblyFormat = [{ $target (`generate_gpu_launch` $generate_gpu_launch^)? (`grid_dims` `=` $grid_dims^)? attr-dict + `:` functional-type($target, $result) }]; let extraClassDeclaration = [{ ::mlir::DiagnosedSilenceableFailure applyToOne( diff --git a/mlir/include/mlir/Dialect/MemRef/TransformOps/MemRefTransformOps.h b/mlir/include/mlir/Dialect/MemRef/TransformOps/MemRefTransformOps.h index dd33df99861c..3d85a1b3a2ce 100644 --- a/mlir/include/mlir/Dialect/MemRef/TransformOps/MemRefTransformOps.h +++ b/mlir/include/mlir/Dialect/MemRef/TransformOps/MemRefTransformOps.h @@ -9,7 +9,6 @@ #ifndef MLIR_DIALECT_MEMREF_TRANSFORMOPS_MEMREFTRANSFORMOPS_H #define MLIR_DIALECT_MEMREF_TRANSFORMOPS_MEMREFTRANSFORMOPS_H -#include "mlir/Dialect/PDL/IR/PDLTypes.h" #include "mlir/Dialect/Transform/IR/TransformInterfaces.h" #include "mlir/IR/OpImplementation.h" diff --git a/mlir/include/mlir/Dialect/MemRef/TransformOps/MemRefTransformOps.td b/mlir/include/mlir/Dialect/MemRef/TransformOps/MemRefTransformOps.td index 1cefe5640833..04dfe1f43024 100644 --- a/mlir/include/mlir/Dialect/MemRef/TransformOps/MemRefTransformOps.td +++ b/mlir/include/mlir/Dialect/MemRef/TransformOps/MemRefTransformOps.td @@ -12,7 +12,6 @@ include "mlir/Dialect/Transform/IR/TransformDialect.td" include "mlir/Dialect/Transform/IR/TransformInterfaces.td" include "mlir/Dialect/Transform/IR/TransformTypes.td" -include "mlir/Dialect/PDL/IR/PDLTypes.td" include "mlir/Interfaces/SideEffectInterfaces.td" include "mlir/IR/OpBase.td" @@ -43,7 +42,7 @@ def MemRefMultiBufferOp : Op<Transform_Dialect, "memref.multibuffer", ConfinedAttr<I64Attr, [IntPositive]>:$factor, UnitAttr:$skip_analysis); - let results = (outs PDL_Operation:$transformed); + let results = (outs TransformHandleTypeInterface:$transformed); let assemblyFormat = "$target attr-dict `:` functional-type(operands, results)"; @@ -81,8 +80,8 @@ def MemRefExtractAddressComputationsOp : to be isolated from above. }]; - let arguments = (ins PDL_Operation:$target); - let results = (outs PDL_Operation:$transformed); + let arguments = (ins TransformHandleTypeInterface:$target); + let results = (outs TransformHandleTypeInterface:$transformed); let assemblyFormat = "$target attr-dict `:` functional-type(operands, results)"; @@ -122,9 +121,10 @@ def MemRefMakeLoopIndependentOp This transform op consumes the target handle and produces a result handle. }]; - let arguments = (ins PDL_Operation:$target, I64Attr:$num_loops); - let results = (outs PDL_Operation:$transformed); - let assemblyFormat = "$target attr-dict"; + let arguments = (ins TransformHandleTypeInterface:$target, I64Attr:$num_loops); + let results = (outs TransformHandleTypeInterface:$transformed); + let assemblyFormat = + "$target attr-dict `:` functional-type($target, $transformed)"; let extraClassDeclaration = [{ ::mlir::DiagnosedSilenceableFailure applyToOne( diff --git a/mlir/include/mlir/Dialect/Tensor/TransformOps/TensorTransformOps.h b/mlir/include/mlir/Dialect/Tensor/TransformOps/TensorTransformOps.h index c735700b7738..e902cce8e7d8 100644 --- a/mlir/include/mlir/Dialect/Tensor/TransformOps/TensorTransformOps.h +++ b/mlir/include/mlir/Dialect/Tensor/TransformOps/TensorTransformOps.h @@ -9,7 +9,6 @@ #ifndef MLIR_DIALECT_TENSOR_TRANSFORMOPS_TENSORTRANSFORMOPS_H #define MLIR_DIALECT_TENSOR_TRANSFORMOPS_TENSORTRANSFORMOPS_H -#include "mlir/Dialect/PDL/IR/PDLTypes.h" #include "mlir/Dialect/Transform/IR/TransformOps.h" #include "mlir/Dialect/Transform/IR/TransformTypes.h" #include "mlir/IR/OpImplementation.h" diff --git a/mlir/include/mlir/Dialect/Tensor/TransformOps/TensorTransformOps.td b/mlir/include/mlir/Dialect/Tensor/TransformOps/TensorTransformOps.td index 42be8821addc..a9580bda3942 100644 --- a/mlir/include/mlir/Dialect/Tensor/TransformOps/TensorTransformOps.td +++ b/mlir/include/mlir/Dialect/Tensor/TransformOps/TensorTransformOps.td @@ -9,7 +9,6 @@ #ifndef TENSOR_TRANSFORM_OPS #define TENSOR_TRANSFORM_OPS -include "mlir/Dialect/PDL/IR/PDLTypes.td" include "mlir/Dialect/Transform/IR/TransformDialect.td" include "mlir/Dialect/Transform/IR/TransformInterfaces.td" include "mlir/Dialect/Transform/IR/TransformTypes.td" @@ -49,9 +48,10 @@ def MakeLoopIndependentOp This transform op consumes the target handle and produces a result handle. }]; - let arguments = (ins PDL_Operation:$target, I64Attr:$num_loops); - let results = (outs PDL_Operation:$transformed); - let assemblyFormat = "$target attr-dict"; + let arguments = (ins TransformHandleTypeInterface:$target, I64Attr:$num_loops); + let results = (outs TransformHandleTypeInterface:$transformed); + let assemblyFormat = + "$target attr-dict `:` functional-type($target, $transformed)"; let extraClassDeclaration = [{ ::mlir::DiagnosedSilenceableFailure applyToOne( diff --git a/mlir/lib/Dialect/Affine/TransformOps/AffineTransformOps.cpp b/mlir/lib/Dialect/Affine/TransformOps/AffineTransformOps.cpp index 9a952e6c9dc5..5a7f092e80bc 100644 --- a/mlir/lib/Dialect/Affine/TransformOps/AffineTransformOps.cpp +++ b/mlir/lib/Dialect/Affine/TransformOps/AffineTransformOps.cpp @@ -12,8 +12,6 @@ #include "mlir/Dialect/Affine/IR/AffineOps.h" #include "mlir/Dialect/Affine/IR/AffineValueMap.h" #include "mlir/Dialect/Affine/LoopUtils.h" -#include "mlir/Dialect/PDL/IR/PDL.h" -#include "mlir/Dialect/PDL/IR/PDLTypes.h" #include "mlir/Dialect/Transform/IR/TransformDialect.h" #include "mlir/Dialect/Transform/IR/TransformInterfaces.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" diff --git a/mlir/lib/Dialect/Affine/TransformOps/CMakeLists.txt b/mlir/lib/Dialect/Affine/TransformOps/CMakeLists.txt index 24e2c8378bb3..4797cd973e8f 100644 --- a/mlir/lib/Dialect/Affine/TransformOps/CMakeLists.txt +++ b/mlir/lib/Dialect/Affine/TransformOps/CMakeLists.txt @@ -12,7 +12,6 @@ add_mlir_dialect_library(MLIRAffineTransformOps MLIRAffineDialect MLIRFuncDialect MLIRIR - MLIRPDLDialect MLIRAffineDialect MLIRAffineTransforms MLIRAffineUtils diff --git a/mlir/lib/Dialect/Bufferization/TransformOps/BufferizationTransformOps.cpp b/mlir/lib/Dialect/Bufferization/TransformOps/BufferizationTransformOps.cpp index 6dd32b815afc..9d45442add5b 100644 --- a/mlir/lib/Dialect/Bufferization/TransformOps/BufferizationTransformOps.cpp +++ b/mlir/lib/Dialect/Bufferization/TransformOps/BufferizationTransformOps.cpp @@ -13,8 +13,6 @@ #include "mlir/Dialect/Bufferization/Transforms/OneShotModuleBufferize.h" #include "mlir/Dialect/Bufferization/Transforms/Transforms.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" -#include "mlir/Dialect/PDL/IR/PDL.h" -#include "mlir/Dialect/PDL/IR/PDLTypes.h" #include "mlir/Dialect/Tensor/IR/Tensor.h" #include "mlir/Dialect/Transform/IR/TransformDialect.h" #include "mlir/IR/FunctionInterfaces.h" @@ -123,8 +121,6 @@ public: using Base::Base; void init() { - declareDependentDialect<pdl::PDLDialect>(); - declareGeneratedDialect<bufferization::BufferizationDialect>(); declareGeneratedDialect<memref::MemRefDialect>(); diff --git a/mlir/lib/Dialect/GPU/TransformOps/CMakeLists.txt b/mlir/lib/Dialect/GPU/TransformOps/CMakeLists.txt index 61695b1d9060..7abd5b93698c 100644 --- a/mlir/lib/Dialect/GPU/TransformOps/CMakeLists.txt +++ b/mlir/lib/Dialect/GPU/TransformOps/CMakeLists.txt @@ -14,7 +14,6 @@ add_mlir_dialect_library(MLIRGPUTransformOps MLIRIR MLIRGPUTransforms MLIRParser - MLIRPDLDialect MLIRSideEffectInterfaces MLIRTransformDialect MLIRGPUDialect diff --git a/mlir/lib/Dialect/GPU/TransformOps/GPUTransformOps.cpp b/mlir/lib/Dialect/GPU/TransformOps/GPUTransformOps.cpp index 8cd2ccfcf882..e9027d103cb6 100644 --- a/mlir/lib/Dialect/GPU/TransformOps/GPUTransformOps.cpp +++ b/mlir/lib/Dialect/GPU/TransformOps/GPUTransformOps.cpp @@ -13,7 +13,6 @@ #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/GPU/IR/GPUDialect.h" #include "mlir/Dialect/GPU/TransformOps/GPUTransformOps.h" -#include "mlir/Dialect/PDL/IR/PDL.h" #include "mlir/Dialect/SCF/IR/DeviceMappingInterface.h" #include "mlir/Dialect/SCF/IR/SCF.h" #include "mlir/Dialect/Transform/IR/TransformDialect.h" @@ -904,7 +903,6 @@ class GPUTransformDialectExtension GPUTransformDialectExtension> { public: GPUTransformDialectExtension() { - declareDependentDialect<pdl::PDLDialect>(); declareGeneratedDialect<scf::SCFDialect>(); declareGeneratedDialect<arith::ArithDialect>(); declareGeneratedDialect<GPUDialect>(); diff --git a/mlir/lib/Dialect/MemRef/TransformOps/CMakeLists.txt b/mlir/lib/Dialect/MemRef/TransformOps/CMakeLists.txt index b32e06aaa09f..0c9ee3d00afe 100644 --- a/mlir/lib/Dialect/MemRef/TransformOps/CMakeLists.txt +++ b/mlir/lib/Dialect/MemRef/TransformOps/CMakeLists.txt @@ -11,7 +11,6 @@ add_mlir_dialect_library(MLIRMemRefTransformOps MLIRAffineDialect MLIRArithDialect MLIRIR - MLIRPDLDialect MLIRLoopLikeInterface MLIRMemRefDialect MLIRMemRefTransforms diff --git a/mlir/lib/Dialect/MemRef/TransformOps/MemRefTransformOps.cpp b/mlir/lib/Dialect/MemRef/TransformOps/MemRefTransformOps.cpp index c523af936e23..f8b449122feb 100644 --- a/mlir/lib/Dialect/MemRef/TransformOps/MemRefTransformOps.cpp +++ b/mlir/lib/Dialect/MemRef/TransformOps/MemRefTransformOps.cpp @@ -13,7 +13,6 @@ #include "mlir/Dialect/MemRef/Transforms/Passes.h" #include "mlir/Dialect/MemRef/Transforms/Transforms.h" #include "mlir/Dialect/NVGPU/IR/NVGPUDialect.h" -#include "mlir/Dialect/PDL/IR/PDL.h" #include "mlir/Dialect/SCF/IR/SCF.h" #include "mlir/Dialect/Transform/IR/TransformDialect.h" #include "mlir/Dialect/Transform/IR/TransformInterfaces.h" @@ -152,7 +151,6 @@ public: using Base::Base; void init() { - declareDependentDialect<pdl::PDLDialect>(); declareGeneratedDialect<affine::AffineDialect>(); declareGeneratedDialect<arith::ArithDialect>(); declareGeneratedDialect<memref::MemRefDialect>(); diff --git a/mlir/lib/Dialect/Tensor/TransformOps/CMakeLists.txt b/mlir/lib/Dialect/Tensor/TransformOps/CMakeLists.txt index be1a5ddf7e7a..27ea4751b45e 100644 --- a/mlir/lib/Dialect/Tensor/TransformOps/CMakeLists.txt +++ b/mlir/lib/Dialect/Tensor/TransformOps/CMakeLists.txt @@ -10,7 +10,6 @@ add_mlir_dialect_library(MLIRTensorTransformOps LINK_LIBS PUBLIC MLIRAffineDialect MLIRIR - MLIRPDLDialect MLIRSCFDialect MLIRTensorDialect MLIRTensorTransforms diff --git a/mlir/test/Dialect/Affine/transform-op-simplify-bounded-affine-ops.mlir b/mlir/test/Dialect/Affine/transform-op-simplify-bounded-affine-ops.mlir index 2f384ea3a251..a30607b07ebe 100644 --- a/mlir/test/Dialect/Affine/transform-op-simplify-bounded-affine-ops.mlir +++ b/mlir/test/Dialect/Affine/transform-op-simplify-bounded-affine-ops.mlir @@ -14,10 +14,10 @@ func.func @simplify_min_max() -> (index, index) { } transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["affine.min", "affine.max"]} in %arg1 : (!pdl.operation) -> !pdl.operation - %1 = transform.structured.match ops{["test.some_op"]} in %arg1 : (!pdl.operation) -> !pdl.operation - transform.affine.simplify_bounded_affine_ops %0 with [%1] within [0] and [20] +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["affine.min", "affine.max"]} in %arg1 : (!transform.any_op) -> !transform.any_op + %1 = transform.structured.match ops{["test.some_op"]} in %arg1 : (!transform.any_op) -> !transform.any_op + transform.affine.simplify_bounded_affine_ops %0 with [%1 : !transform.any_op] within [0] and [20] : !transform.any_op } // ----- @@ -34,27 +34,27 @@ func.func @simplify_min_sequence() -> index { } transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["affine.min"]} in %arg1 : (!pdl.operation) -> !pdl.operation - %1 = transform.structured.match ops{["test.workgroup_id"]} in %arg1 : (!pdl.operation) -> !pdl.operation - %2 = transform.structured.match ops{["test.thread_id"]} in %arg1 : (!pdl.operation) -> !pdl.operation - transform.affine.simplify_bounded_affine_ops %0 with [%1, %2] within [0, 0] and [31, 31] +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["affine.min"]} in %arg1 : (!transform.any_op) -> !transform.any_op + %1 = transform.structured.match ops{["test.workgroup_id"]} in %arg1 : (!transform.any_op) -> !transform.any_op + %2 = transform.structured.match ops{["test.thread_id"]} in %arg1 : (!transform.any_op) -> !transform.any_op + transform.affine.simplify_bounded_affine_ops %0 with [%1, %2 : !transform.any_op, !transform.any_op] within [0, 0] and [31, 31] : !transform.any_op } // ----- transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["affine.min"]} in %arg1 : (!pdl.operation) -> !pdl.operation +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["affine.min"]} in %arg1 : (!transform.any_op) -> !transform.any_op // expected-error@+1 {{incorrect number of lower bounds, expected 0 but found 1}} - transform.affine.simplify_bounded_affine_ops %0 with [] within [0] and [] + transform.affine.simplify_bounded_affine_ops %0 with [] within [0] and [] : !transform.any_op } // ----- transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["affine.min"]} in %arg1 : (!pdl.operation) -> !pdl.operation +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["affine.min"]} in %arg1 : (!transform.any_op) -> !transform.any_op // expected-error@+1 {{incorrect number of upper bounds, expected 0 but found 1}} - transform.affine.simplify_bounded_affine_ops %0 with [] within [] and [5] + transform.affine.simplify_bounded_affine_ops %0 with [] within [] and [5] : !transform.any_op } diff --git a/mlir/test/Dialect/Bufferization/Transforms/transform-ops.mlir b/mlir/test/Dialect/Bufferization/Transforms/transform-ops.mlir index 05d6a1bbdd55..c4a404489194 100644 --- a/mlir/test/Dialect/Bufferization/Transforms/transform-ops.mlir +++ b/mlir/test/Dialect/Bufferization/Transforms/transform-ops.mlir @@ -3,9 +3,9 @@ // Test One-Shot Bufferize. transform.sequence failures(propagate) { -^bb0(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!pdl.operation) -> !pdl.operation - %1 = transform.bufferization.one_shot_bufferize %0 : (!pdl.operation) -> !pdl.operation +^bb0(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!transform.any_op) -> !transform.any_op + %1 = transform.bufferization.one_shot_bufferize %0 : (!transform.any_op) -> !transform.any_op } // CHECK-LABEL: func @test_function( @@ -31,10 +31,10 @@ func.func @test_function(%A : tensor<?xf32>, %v : vector<4xf32>) -> (tensor<?xf3 // Test analysis of One-Shot Bufferize only. transform.sequence failures(propagate) { -^bb0(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!pdl.operation) -> !pdl.operation +^bb0(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!transform.any_op) -> !transform.any_op %1 = transform.bufferization.one_shot_bufferize %0 - {test_analysis_only = true} : (!pdl.operation) -> !pdl.operation + {test_analysis_only = true} : (!transform.any_op) -> !transform.any_op } // CHECK-LABEL: func @test_function_analysis( @@ -54,10 +54,10 @@ func.func @test_function_analysis(%A : tensor<?xf32>, %v : vector<4xf32>) -> (te // allowed with `allow_unknown_ops`. transform.sequence failures(propagate) { -^bb0(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!pdl.operation) -> !pdl.operation +^bb0(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!transform.any_op) -> !transform.any_op // expected-error @+1 {{bufferization failed}} - %1 = transform.bufferization.one_shot_bufferize %0 : (!pdl.operation) -> !pdl.operation + %1 = transform.bufferization.one_shot_bufferize %0 : (!transform.any_op) -> !transform.any_op } func.func @test_unknown_op_failure() -> (tensor<?xf32>) { @@ -69,9 +69,9 @@ func.func @test_unknown_op_failure() -> (tensor<?xf32>) { // ----- transform.sequence failures(propagate) { -^bb0(%arg1: !pdl.operation): +^bb0(%arg1: !transform.any_op): // %arg1 is the module - %0 = transform.bufferization.one_shot_bufferize %arg1 : (!pdl.operation) -> !pdl.operation + %0 = transform.bufferization.one_shot_bufferize %arg1 : (!transform.any_op) -> !transform.any_op } module { @@ -99,9 +99,9 @@ module { // Test we use identity layout at function boundaries. transform.sequence failures(propagate) { - ^bb0(%arg1: !pdl.operation): + ^bb0(%arg1: !transform.any_op): %0 = transform.bufferization.one_shot_bufferize layout{IdentityLayoutMap} %arg1 - { bufferize_function_boundaries = true } : (!pdl.operation) -> !pdl.operation + { bufferize_function_boundaries = true } : (!transform.any_op) -> !transform.any_op } // CHECK: func.func @matmul( @@ -118,9 +118,9 @@ func.func @matmul(%A: tensor<12x9xf32>, %B: tensor<9x6xf32>, %C: tensor<12x6xf32 // ----- transform.sequence failures(propagate) { - ^bb0(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["tensor.empty"]} in %arg1 : (!pdl.operation) -> !pdl.operation - %1 = transform.cast %0 : !pdl.operation to !transform.op<"tensor.empty"> + ^bb0(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["tensor.empty"]} in %arg1 : (!transform.any_op) -> !transform.any_op + %1 = transform.cast %0 : !transform.any_op to !transform.op<"tensor.empty"> transform.bufferization.empty_tensor_to_alloc_tensor %1 : (!transform.op<"tensor.empty">) -> !transform.op<"bufferization.alloc_tensor"> } @@ -134,9 +134,9 @@ func.func @empty_to_tensor_alloc() -> tensor<2x2xf32> { // ----- transform.sequence failures(propagate) { -^bb0(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!pdl.operation) -> !pdl.operation - transform.bufferization.eliminate_empty_tensors %0 +^bb0(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!transform.any_op) -> !transform.any_op + transform.bufferization.eliminate_empty_tensors %0 : !transform.any_op } // CHECK-LABEL: func @empty_tensor_elimination( diff --git a/mlir/test/Dialect/GPU/transform-gpu-failing.mlir b/mlir/test/Dialect/GPU/transform-gpu-failing.mlir index 813b6f3cd17e..0d560482d651 100644 --- a/mlir/test/Dialect/GPU/transform-gpu-failing.mlir +++ b/mlir/test/Dialect/GPU/transform-gpu-failing.mlir @@ -5,10 +5,10 @@ func.func @map_nested_forall_to_threads_not_gpu_launch() -> () { return } transform.sequence failures(propagate) { -^bb0(%arg0: !pdl.operation): - %funcop = transform.structured.match ops{["tensor.empty"]} in %arg0 : (!pdl.operation) -> !pdl.operation +^bb0(%arg0: !transform.any_op): + %funcop = transform.structured.match ops{["tensor.empty"]} in %arg0 : (!transform.any_op) -> !transform.any_op // expected-error @below {{Given target is not a gpu.launch}} - %1 = transform.gpu.map_nested_forall_to_threads %funcop block_dims = [1, 1, 1] + %1 = transform.gpu.map_nested_forall_to_threads %funcop block_dims = [1, 1, 1] : (!transform.any_op) -> !transform.any_op } // ----- @@ -45,11 +45,11 @@ func.func @map_nested_forall_to_threads_excessive_threads(%x: memref<2 x 32 x f3 return %y : memref<2 x 32 x f32> } transform.sequence failures(propagate) { -^bb1(%arg0: !pdl.operation): - %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!pdl.operation) -> !pdl.operation +^bb1(%arg0: !transform.any_op): + %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!transform.any_op) -> !transform.any_op // expected-error @below {{Trying to launch a GPU kernel with grid_dims = (1, 1, 1) block_dims = (1200, 9, 1). It is larger than the limits.}} // expected-note @below {{"block_dims" is too large}} - transform.gpu.map_nested_forall_to_threads %funcop block_dims = [1200, 9, 1] + transform.gpu.map_nested_forall_to_threads %funcop block_dims = [1200, 9, 1] : (!transform.any_op) -> !transform.any_op } // ----- @@ -87,10 +87,10 @@ func.func @map_nested_forall_to_threads_fewer_threads(%x: memref<2 x 32 x f32>, } transform.sequence failures(propagate) { -^bb1(%arg0: !pdl.operation): - %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!pdl.operation) -> !pdl.operation +^bb1(%arg0: !transform.any_op): + %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!transform.any_op) -> !transform.any_op // expected-error @below {{Trying to map to fewer GPU threads than loop iterations but overprovisioning is not yet supported. Try additional tiling of the before mapping or map to more threads.}} - transform.gpu.map_nested_forall_to_threads %funcop block_dims = [128, 4, 1] + transform.gpu.map_nested_forall_to_threads %funcop block_dims = [128, 4, 1] : (!transform.any_op) -> !transform.any_op } // ----- @@ -113,10 +113,10 @@ func.func @map_nested_forall_to_threads_dynamic_trip_count(%x: memref<2 x 32 x f } transform.sequence failures(propagate) { -^bb1(%arg0: !pdl.operation): - %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!pdl.operation) -> !pdl.operation +^bb1(%arg0: !transform.any_op): + %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!transform.any_op) -> !transform.any_op // expected-error @below {{unsupported dynamic sizes}} - transform.gpu.map_nested_forall_to_threads %funcop block_dims = [128, 4, 1] + transform.gpu.map_nested_forall_to_threads %funcop block_dims = [128, 4, 1] : (!transform.any_op) -> !transform.any_op } // ----- @@ -137,9 +137,9 @@ transform.sequence failures(propagate) { %matmul = transform.structured.match ops{["linalg.matmul"]} in %arg0 : (!transform.any_op) -> !transform.any_op %forall, %tiled = transform.structured.tile_to_forall_op %matmul num_threads [10, 20, 30] (mapping = [ #gpu.thread<y>, #gpu.thread<x>, #gpu.thread<z> ] ) : (!transform.any_op) -> (!transform.any_op, !transform.any_op) - %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!transform.any_op) -> !pdl.operation + %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!transform.any_op) -> !transform.any_op // expected-error @below {{only bufferized scf.forall can be mapped}} - transform.gpu.map_nested_forall_to_threads %funcop block_dims = [128, 4, 1] + transform.gpu.map_nested_forall_to_threads %funcop block_dims = [128, 4, 1] : (!transform.any_op) -> !transform.any_op } // ----- @@ -151,10 +151,10 @@ func.func @map_forall_to_blocks_not_gpu_launch() -> () { return } transform.sequence failures(propagate) { -^bb0(%arg0: !pdl.operation): - %funcop = transform.structured.match ops{["tensor.empty"]} in %arg0 : (!pdl.operation) -> !pdl.operation +^bb0(%arg0: !transform.any_op): + %funcop = transform.structured.match ops{["tensor.empty"]} in %arg0 : (!transform.any_op) -> !transform.any_op // expected-error @below {{Given target is not gpu.launch}} - %1 = transform.gpu.map_forall_to_blocks %funcop + %1 = transform.gpu.map_forall_to_blocks %funcop : (!transform.any_op) -> !transform.any_op } // ----- @@ -188,10 +188,10 @@ func.func @map_forall_to_blocks_not_unique(%x: memref<2 x 32 x f32>, %y: memref< } transform.sequence failures(propagate) { -^bb0(%arg0: !pdl.operation): - %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!pdl.operation) -> !pdl.operation +^bb0(%arg0: !transform.any_op): + %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!transform.any_op) -> !transform.any_op // expected-error @below {{could not find a unique topLevel scf.forall}} - %1 = transform.gpu.map_forall_to_blocks %funcop + %1 = transform.gpu.map_forall_to_blocks %funcop : (!transform.any_op) -> !transform.any_op } // ----- @@ -221,10 +221,10 @@ func.func @map_forall_to_blocks_large_loop(%x: memref<2 x 32 x f32>, %y: memref< } transform.sequence failures(propagate) { -^bb0(%arg0: !pdl.operation): - %funcop = transform.structured.match ops{["func.func"]} in %arg0 : (!pdl.operation) -> !pdl.operation +^bb0(%arg0: !transform.any_op): + %funcop = transform.structured.match ops{["func.func"]} in %arg0 : (!transform.any_op) -> !transform.any_op // expected-error @below {{could not find a unique topLevel scf.forall}} - %1 = transform.gpu.map_forall_to_blocks %funcop { generate_gpu_launch } + %1 = transform.gpu.map_forall_to_blocks %funcop { generate_gpu_launch } : (!transform.any_op) -> !transform.any_op } // ----- @@ -242,10 +242,10 @@ func.func @map_forall_to_blocks_large_loop(%x: memref<2 x 32 x f32>, %y: memref< } transform.sequence failures(propagate) { -^bb0(%arg0: !pdl.operation): - %funcop = transform.structured.match ops{["func.func"]} in %arg0 : (!pdl.operation) -> !pdl.operation +^bb0(%arg0: !transform.any_op): + %funcop = transform.structured.match ops{["func.func"]} in %arg0 : (!transform.any_op) -> !transform.any_op // expected-error @below {{Trying to launch a GPU kernel with grid_dims = (65535, 65535, 1) block_dims = (1, 1, 1). It is larger than the limits.}} - %1 = transform.gpu.map_forall_to_blocks %funcop generate_gpu_launch + %1 = transform.gpu.map_forall_to_blocks %funcop generate_gpu_launch : (!transform.any_op) -> !transform.any_op } // ----- @@ -269,10 +269,10 @@ func.func @saxpy2d_singleloop(%x: !type, %y: !type, %stream : !gpu.async.token) } transform.sequence failures(propagate) { -^bb1(%arg0: !pdl.operation): - %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!pdl.operation) -> !pdl.operation +^bb1(%arg0: !transform.any_op): + %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!transform.any_op) -> !transform.any_op // expected-error @below {{duplicated attribute, cannot map different loops to the same processor}} - transform.gpu.map_nested_forall_to_threads %funcop block_dims = [32, 32, 1] + transform.gpu.map_nested_forall_to_threads %funcop block_dims = [32, 32, 1] : (!transform.any_op) -> !transform.any_op } // ----- diff --git a/mlir/test/Dialect/GPU/transform-gpu.mlir b/mlir/test/Dialect/GPU/transform-gpu.mlir index d9872a9666f6..64d80f95a6b3 100644 --- a/mlir/test/Dialect/GPU/transform-gpu.mlir +++ b/mlir/test/Dialect/GPU/transform-gpu.mlir @@ -31,9 +31,9 @@ func.func @saxpy2dblock(%x: !type, %y: !type, %t: !type1d, %alpha : f32, %stream } transform.sequence failures(propagate) { -^bb1(%arg0: !pdl.operation): - %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!pdl.operation) -> !pdl.operation - transform.gpu.map_forall_to_blocks %funcop grid_dims = [12, 9, 1] +^bb1(%arg0: !transform.any_op): + %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!transform.any_op) -> !transform.any_op + transform.gpu.map_forall_to_blocks %funcop grid_dims = [12, 9, 1] : (!transform.any_op) -> !transform.any_op } // ----- @@ -85,9 +85,9 @@ func.func @saxpy2d(%x: !type, %y: !type, %t: !type1d, %alpha : f32, %stream : !g } transform.sequence failures(propagate) { -^bb1(%arg0: !pdl.operation): - %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!pdl.operation) -> !pdl.operation - transform.gpu.map_nested_forall_to_threads %funcop block_dims = [12, 9, 1] +^bb1(%arg0: !transform.any_op): + %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!transform.any_op) -> !transform.any_op + transform.gpu.map_nested_forall_to_threads %funcop block_dims = [12, 9, 1] : (!transform.any_op) -> !transform.any_op } // ----- @@ -124,10 +124,10 @@ func.func @saxpy4d(%x: !type4d, %y: !type4d, %alpha : f32) -> !type4d { } transform.sequence failures(propagate) { -^bb1(%arg0: !pdl.operation): - %funcop = transform.structured.match ops{["func.func"]} in %arg0 : (!pdl.operation) -> !pdl.operation - %gpuLaunch = transform.gpu.map_forall_to_blocks %funcop { generate_gpu_launch } - transform.gpu.map_nested_forall_to_threads %gpuLaunch block_dims = [32, 4, 1] +^bb1(%arg0: !transform.any_op): + %funcop = transform.structured.match ops{["func.func"]} in %arg0 : (!transform.any_op) -> !transform.any_op + %gpuLaunch = transform.gpu.map_forall_to_blocks %funcop { generate_gpu_launch } : (!transform.any_op) -> !transform.any_op + transform.gpu.map_nested_forall_to_threads %gpuLaunch block_dims = [32, 4, 1] : (!transform.any_op) -> !transform.any_op } // ----- @@ -158,9 +158,9 @@ func.func @saxpy2d_no_barrier(%x: !type, %y: !type, %t: !type1d, %alpha : f32, % } transform.sequence failures(propagate) { -^bb1(%arg0: !pdl.operation): - %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!pdl.operation) -> !pdl.operation - transform.gpu.map_nested_forall_to_threads %funcop block_dims = [12, 9, 1] sync_after_distribute = false +^bb1(%arg0: !transform.any_op): + %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!transform.any_op) -> !transform.any_op + transform.gpu.map_nested_forall_to_threads %funcop block_dims = [12, 9, 1] sync_after_distribute = false : (!transform.any_op) -> !transform.any_op } // ----- @@ -190,9 +190,9 @@ func.func @saxpy2d_singleloop(%x: !type, %y: !type, %stream : !gpu.async.token) } transform.sequence failures(propagate) { -^bb1(%arg0: !pdl.operation): - %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!pdl.operation) -> !pdl.operation - transform.gpu.map_nested_forall_to_threads %funcop block_dims = [32, 1, 1] +^bb1(%arg0: !transform.any_op): + %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!transform.any_op) -> !transform.any_op + transform.gpu.map_nested_forall_to_threads %funcop block_dims = [32, 1, 1] : (!transform.any_op) -> !transform.any_op } // ----- @@ -226,9 +226,9 @@ func.func @saxpy3d_fold_id_z(%x: !type, %y: !type, %t: !type1d, %alpha : f32, %s } transform.sequence failures(propagate) { -^bb1(%arg0: !pdl.operation): - %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!pdl.operation) -> !pdl.operation - transform.gpu.map_nested_forall_to_threads %funcop block_dims = [12, 9, 1] sync_after_distribute = false +^bb1(%arg0: !transform.any_op): + %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!transform.any_op) -> !transform.any_op + transform.gpu.map_nested_forall_to_threads %funcop block_dims = [12, 9, 1] sync_after_distribute = false : (!transform.any_op) -> !transform.any_op } // ----- @@ -302,8 +302,8 @@ func.func @map_multi_level(%x: !type, %y: !type, %t: !type1d, %alpha : f32, %str } transform.sequence failures(propagate) { -^bb1(%arg0: !pdl.operation): - %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!pdl.operation) -> !pdl.operation +^bb1(%arg0: !transform.any_op): + %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0 : (!transform.any_op) -> !transform.any_op transform.gpu.map_nested_forall_to_threads %funcop - block_dims = [12, 11, 1] warp_dims = [3, 2, 1] + block_dims = [12, 11, 1] warp_dims = [3, 2, 1] : (!transform.any_op) -> !transform.any_op } diff --git a/mlir/test/Dialect/MemRef/extract-address-computations.mlir b/mlir/test/Dialect/MemRef/extract-address-computations.mlir index 17e2ac3bc5e2..5064f60d2f7b 100644 --- a/mlir/test/Dialect/MemRef/extract-address-computations.mlir +++ b/mlir/test/Dialect/MemRef/extract-address-computations.mlir @@ -22,11 +22,11 @@ func.func @test_load(%base : memref<2x16x16xf32>, %offset : index) -> f32 { } transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!pdl.operation) -> !pdl.operation - %1 = transform.memref.extract_address_computations %0 : (!pdl.operation) -> !pdl.operation +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!transform.any_op) -> !transform.any_op + %1 = transform.memref.extract_address_computations %0 : (!transform.any_op) -> !transform.any_op // Verify that the returned handle is usable. - transform.test_print_remark_at_operand %1, "transformed" : !pdl.operation + transform.test_print_remark_at_operand %1, "transformed" : !transform.any_op } // ----- @@ -48,9 +48,9 @@ func.func @test_load_nontemporal(%base : memref<2x16x16xf32>, %offset : index) - } transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!pdl.operation) -> !pdl.operation - %1 = transform.memref.extract_address_computations %0 : (!pdl.operation) -> !pdl.operation +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!transform.any_op) -> !transform.any_op + %1 = transform.memref.extract_address_computations %0 : (!transform.any_op) -> !transform.any_op } // ----- @@ -77,9 +77,9 @@ func.func @test_store(%base : memref<2x16x16xf32>, %offset : index) -> () { } transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!pdl.operation) -> !pdl.operation - %1 = transform.memref.extract_address_computations %0 : (!pdl.operation) -> !pdl.operation +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!transform.any_op) -> !transform.any_op + %1 = transform.memref.extract_address_computations %0 : (!transform.any_op) -> !transform.any_op } // ----- @@ -103,9 +103,9 @@ func.func @test_store_nontemporal(%base : memref<2x16x16xf32>, %offset : index) } transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!pdl.operation) -> !pdl.operation - %1 = transform.memref.extract_address_computations %0 : (!pdl.operation) -> !pdl.operation +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!transform.any_op) -> !transform.any_op + %1 = transform.memref.extract_address_computations %0 : (!transform.any_op) -> !transform.any_op } // ----- @@ -157,9 +157,9 @@ func.func @testWithLoop(%base : memref<?x?x?xf32, strided<[?,?,?], offset: ?>>) } transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!pdl.operation) -> !pdl.operation - %1 = transform.memref.extract_address_computations %0 : (!pdl.operation) -> !pdl.operation +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!transform.any_op) -> !transform.any_op + %1 = transform.memref.extract_address_computations %0 : (!transform.any_op) -> !transform.any_op } // ----- @@ -195,9 +195,9 @@ func.func @test_ldmatrix(%base : memref<4x32x32xf16, 3>, } transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!pdl.operation) -> !pdl.operation - %1 = transform.memref.extract_address_computations %0 : (!pdl.operation) -> !pdl.operation +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!transform.any_op) -> !transform.any_op + %1 = transform.memref.extract_address_computations %0 : (!transform.any_op) -> !transform.any_op } // ----- @@ -229,9 +229,9 @@ func.func @test_ldmatrix(%base : memref<?x?x?xf16, 3>, } transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!pdl.operation) -> !pdl.operation - %1 = transform.memref.extract_address_computations %0 : (!pdl.operation) -> !pdl.operation +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!transform.any_op) -> !transform.any_op + %1 = transform.memref.extract_address_computations %0 : (!transform.any_op) -> !transform.any_op } // ----- @@ -264,9 +264,9 @@ func.func @test_transfer_read_op(%base : memref<?x?x?xf16>, } transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!pdl.operation) -> !pdl.operation - %1 = transform.memref.extract_address_computations %0 : (!pdl.operation) -> !pdl.operation +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!transform.any_op) -> !transform.any_op + %1 = transform.memref.extract_address_computations %0 : (!transform.any_op) -> !transform.any_op } // ----- @@ -292,9 +292,9 @@ func.func @test_transfer_read_op_with_tensor(%base : tensor<?x?x?xf16>, } transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!pdl.operation) -> !pdl.operation - %1 = transform.memref.extract_address_computations %0 : (!pdl.operation) -> !pdl.operation +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!transform.any_op) -> !transform.any_op + %1 = transform.memref.extract_address_computations %0 : (!transform.any_op) -> !transform.any_op } // ----- @@ -326,9 +326,9 @@ func.func @test_transfer_write_op(%base : memref<?x?x?xf16>, } transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!pdl.operation) -> !pdl.operation - %1 = transform.memref.extract_address_computations %0 : (!pdl.operation) -> !pdl.operation +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!transform.any_op) -> !transform.any_op + %1 = transform.memref.extract_address_computations %0 : (!transform.any_op) -> !transform.any_op } // ----- @@ -361,9 +361,9 @@ func.func @test_transfer_write_op_with_strides(%base : memref<?x?x?xf16, strided } transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!pdl.operation) -> !pdl.operation - %1 = transform.memref.extract_address_computations %0 : (!pdl.operation) -> !pdl.operation +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!transform.any_op) -> !transform.any_op + %1 = transform.memref.extract_address_computations %0 : (!transform.any_op) -> !transform.any_op } // ----- @@ -387,7 +387,7 @@ func.func @test_transfer_write_op_with_tensor(%base : tensor<?x?x?xf16>, } transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!pdl.operation) -> !pdl.operation - %1 = transform.memref.extract_address_computations %0 : (!pdl.operation) -> !pdl.operation +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["func.func"]} in %arg1 : (!transform.any_op) -> !transform.any_op + %1 = transform.memref.extract_address_computations %0 : (!transform.any_op) -> !transform.any_op } diff --git a/mlir/test/Dialect/MemRef/make-loop-independent.mlir b/mlir/test/Dialect/MemRef/make-loop-independent.mlir index 86ce9f4df400..1a34d9ce74f2 100644 --- a/mlir/test/Dialect/MemRef/make-loop-independent.mlir +++ b/mlir/test/Dialect/MemRef/make-loop-independent.mlir @@ -36,9 +36,9 @@ func.func @make_alloca_loop_independent(%lb: index, %ub: index, %step: index) { return } transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["memref.alloca"]} in %arg1 : (!pdl.operation) -> !pdl.operation - %1 = transform.memref.make_loop_independent %0 {num_loops = 1} +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["memref.alloca"]} in %arg1 : (!transform.any_op) -> !transform.any_op + %1 = transform.memref.make_loop_independent %0 {num_loops = 1} : (!transform.any_op) -> !transform.any_op } // ----- @@ -68,7 +68,7 @@ func.func @make_alloca_loop_independent_static(%step: index) { return } transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["memref.alloca"]} in %arg1 : (!pdl.operation) -> !pdl.operation - %1 = transform.memref.make_loop_independent %0 {num_loops = 1} +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["memref.alloca"]} in %arg1 : (!transform.any_op) -> !transform.any_op + %1 = transform.memref.make_loop_independent %0 {num_loops = 1} : (!transform.any_op) -> !transform.any_op } diff --git a/mlir/test/Dialect/MemRef/transform-ops.mlir b/mlir/test/Dialect/MemRef/transform-ops.mlir index 3edcc53230fa..8d2d56921621 100644 --- a/mlir/test/Dialect/MemRef/transform-ops.mlir +++ b/mlir/test/Dialect/MemRef/transform-ops.mlir @@ -29,11 +29,11 @@ func.func @multi_buffer(%in: memref<16xf32>) { } transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["memref.alloc"]} in %arg1 : (!pdl.operation) -> !transform.op<"memref.alloc"> - %1 = transform.memref.multibuffer %0 {factor = 2 : i64} : (!transform.op<"memref.alloc">) -> !pdl.operation +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["memref.alloc"]} in %arg1 : (!transform.any_op) -> !transform.op<"memref.alloc"> + %1 = transform.memref.multibuffer %0 {factor = 2 : i64} : (!transform.op<"memref.alloc">) -> !transform.any_op // Verify that the returned handle is usable. - transform.test_print_remark_at_operand %1, "transformed" : !pdl.operation + transform.test_print_remark_at_operand %1, "transformed" : !transform.any_op } // ----- @@ -64,11 +64,11 @@ func.func @multi_buffer_on_affine_loop(%in: memref<16xf32>) { } transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["memref.alloc"]} in %arg1 : (!pdl.operation) -> !transform.op<"memref.alloc"> - %1 = transform.memref.multibuffer %0 {factor = 2 : i64} : (!transform.op<"memref.alloc">) -> !pdl.operation +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["memref.alloc"]} in %arg1 : (!transform.any_op) -> !transform.op<"memref.alloc"> + %1 = transform.memref.multibuffer %0 {factor = 2 : i64} : (!transform.op<"memref.alloc">) -> !transform.any_op // Verify that the returned handle is usable. - transform.test_print_remark_at_operand %1, "transformed" : !pdl.operation + transform.test_print_remark_at_operand %1, "transformed" : !transform.any_op } // ----- @@ -102,9 +102,9 @@ func.func @multi_buffer_uses_with_no_loop_dominator(%in: memref<16xf32>, %cond: } transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["memref.alloc"]} in %arg1 : (!pdl.operation) -> !transform.op<"memref.alloc"> - %1 = transform.memref.multibuffer %0 {factor = 2 : i64} : (!transform.op<"memref.alloc">) -> !pdl.operation +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["memref.alloc"]} in %arg1 : (!transform.any_op) -> !transform.op<"memref.alloc"> + %1 = transform.memref.multibuffer %0 {factor = 2 : i64} : (!transform.op<"memref.alloc">) -> !transform.any_op } // ----- @@ -137,10 +137,10 @@ func.func @multi_buffer_reject_alloca(%in: memref<16xf32>, %cond: i1) { } transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["memref.alloca"]} in %arg1 : (!pdl.operation) -> !transform.op<"memref.alloca"> +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["memref.alloca"]} in %arg1 : (!transform.any_op) -> !transform.op<"memref.alloca"> // expected-error @below {{'transform.memref.multibuffer' op operand #0 must be Transform IR handle to memref.alloc operations, but got '!transform.op<"memref.alloca">'}} - %1 = transform.memref.multibuffer %0 {factor = 2 : i64} : (!transform.op<"memref.alloca">) -> !pdl.operation + %1 = transform.memref.multibuffer %0 {factor = 2 : i64} : (!transform.op<"memref.alloca">) -> !transform.any_op } // ----- @@ -179,11 +179,11 @@ func.func @multi_buffer_one_alloc_with_use_outside_of_loop(%in: memref<16xf32>) } transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["memref.alloc"]} in %arg1 : (!pdl.operation) -> !transform.op<"memref.alloc"> - %1 = transform.memref.multibuffer %0 {factor = 2 : i64} : (!transform.op<"memref.alloc">) -> !pdl.operation +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["memref.alloc"]} in %arg1 : (!transform.any_op) -> !transform.op<"memref.alloc"> + %1 = transform.memref.multibuffer %0 {factor = 2 : i64} : (!transform.op<"memref.alloc">) -> !transform.any_op // Verify that the returned handle is usable. - transform.test_print_remark_at_operand %1, "transformed" : !pdl.operation + transform.test_print_remark_at_operand %1, "transformed" : !transform.any_op } // ----- @@ -213,11 +213,11 @@ func.func @multi_buffer_no_analysis(%in: memref<16xf32>) { } transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["memref.alloc"]} in %arg1 : (!pdl.operation) -> !transform.op<"memref.alloc"> - %1 = transform.memref.multibuffer %0 {factor = 2 : i64, skip_analysis} : (!transform.op<"memref.alloc">) -> !pdl.operation +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["memref.alloc"]} in %arg1 : (!transform.any_op) -> !transform.op<"memref.alloc"> + %1 = transform.memref.multibuffer %0 {factor = 2 : i64, skip_analysis} : (!transform.op<"memref.alloc">) -> !transform.any_op // Verify that the returned handle is usable. - transform.test_print_remark_at_operand %1, "transformed" : !pdl.operation + transform.test_print_remark_at_operand %1, "transformed" : !transform.any_op } // ----- @@ -250,9 +250,9 @@ func.func @multi_buffer_dealloc(%in: memref<16xf32>) { } transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["memref.alloc"]} in %arg1 : (!pdl.operation) -> !transform.op<"memref.alloc"> - %1 = transform.memref.multibuffer %0 {factor = 2 : i64, skip_analysis} : (!transform.op<"memref.alloc">) -> !pdl.operation +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["memref.alloc"]} in %arg1 : (!transform.any_op) -> !transform.op<"memref.alloc"> + %1 = transform.memref.multibuffer %0 {factor = 2 : i64, skip_analysis} : (!transform.op<"memref.alloc">) -> !transform.any_op // Verify that the returned handle is usable. - transform.test_print_remark_at_operand %1, "transformed" : !pdl.operation + transform.test_print_remark_at_operand %1, "transformed" : !transform.any_op } diff --git a/mlir/test/Dialect/Tensor/transform-op-make-loop-independent.mlir b/mlir/test/Dialect/Tensor/transform-op-make-loop-independent.mlir index 18a99c5c437b..d379888c1468 100644 --- a/mlir/test/Dialect/Tensor/transform-op-make-loop-independent.mlir +++ b/mlir/test/Dialect/Tensor/transform-op-make-loop-independent.mlir @@ -30,9 +30,9 @@ func.func @make_pad_loop_independent_1(%lb: index, %ub: index, %step: index, } transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["tensor.pad"]} in %arg1 : (!pdl.operation) -> !pdl.operation - %1 = transform.tensor.make_loop_independent %0 {num_loops = 1} +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["tensor.pad"]} in %arg1 : (!transform.any_op) -> !transform.any_op + %1 = transform.tensor.make_loop_independent %0 {num_loops = 1} : (!transform.any_op) -> !transform.any_op } // ----- @@ -67,9 +67,9 @@ func.func @make_pad_loop_independent_1(%lb: index, %ub: index, %step: index, } transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["tensor.pad"]} in %arg1 : (!pdl.operation) -> !pdl.operation - %1 = transform.tensor.make_loop_independent %0 {num_loops = 1} +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["tensor.pad"]} in %arg1 : (!transform.any_op) -> !transform.any_op + %1 = transform.tensor.make_loop_independent %0 {num_loops = 1} : (!transform.any_op) -> !transform.any_op } // ----- @@ -93,9 +93,9 @@ func.func @two_loops(%lb: index, %ub: index, %step: index, } transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["tensor.pad"]} in %arg1 : (!pdl.operation) -> !pdl.operation - %1 = transform.tensor.make_loop_independent %0 {num_loops = 2} +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["tensor.pad"]} in %arg1 : (!transform.any_op) -> !transform.any_op + %1 = transform.tensor.make_loop_independent %0 {num_loops = 2} : (!transform.any_op) -> !transform.any_op } // ----- @@ -117,10 +117,10 @@ func.func @not_enough_loops(%lb: index, %ub: index, %step: index, } transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["tensor.pad"]} in %arg1 : (!pdl.operation) -> !pdl.operation +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["tensor.pad"]} in %arg1 : (!transform.any_op) -> !transform.any_op // expected-error@below {{could not find 2-th enclosing loop}} - %1 = transform.tensor.make_loop_independent %0 {num_loops = 3} + %1 = transform.tensor.make_loop_independent %0 {num_loops = 3} : (!transform.any_op) -> !transform.any_op } // ----- @@ -145,7 +145,7 @@ func.func @make_empty_loop_independent(%lb: index, %ub: index, %step: index) { } transform.sequence failures(propagate) { -^bb1(%arg1: !pdl.operation): - %0 = transform.structured.match ops{["tensor.empty"]} in %arg1 : (!pdl.operation) -> !pdl.operation - %1 = transform.tensor.make_loop_independent %0 {num_loops = 1} +^bb1(%arg1: !transform.any_op): + %0 = transform.structured.match ops{["tensor.empty"]} in %arg1 : (!transform.any_op) -> !transform.any_op + %1 = transform.tensor.make_loop_independent %0 {num_loops = 1} : (!transform.any_op) -> !transform.any_op } |