From 2f3ac28cb2f7fc24c6ff742af571b58419c0adaa Mon Sep 17 00:00:00 2001 From: Alex Zinenko Date: Wed, 17 May 2023 14:01:08 +0000 Subject: [PATCH] [mlir] don't hardcode PDL_Operation in Transform dialect extensions Update operations in Transform dialect extensions defined in the Affine, GPU, MemRef and Tensor dialects to use the more generic `TransformHandleTypeInterface` type constraint instead of hardcoding `PDL_Operation`. See https://discourse.llvm.org/t/rfc-type-system-for-the-transform-dialect/65702 for motivation. Remove the dependency on PDLDialect from these extensions. Update tests to use `!transform.any_op` instead of `!pdl.operation`. Reviewed By: nicolasvasilache Differential Revision: https://reviews.llvm.org/D150781 --- .../Affine/TransformOps/AffineTransformOps.h | 1 - .../Affine/TransformOps/AffineTransformOps.td | 8 +-- .../TransformOps/BufferizationTransformOps.h | 1 - .../TransformOps/BufferizationTransformOps.td | 5 +- .../Dialect/GPU/TransformOps/GPUTransformOps.h | 1 - .../Dialect/GPU/TransformOps/GPUTransformOps.td | 13 ++-- .../MemRef/TransformOps/MemRefTransformOps.h | 1 - .../MemRef/TransformOps/MemRefTransformOps.td | 14 ++-- .../Tensor/TransformOps/TensorTransformOps.h | 1 - .../Tensor/TransformOps/TensorTransformOps.td | 8 +-- .../Affine/TransformOps/AffineTransformOps.cpp | 2 - .../lib/Dialect/Affine/TransformOps/CMakeLists.txt | 1 - .../TransformOps/BufferizationTransformOps.cpp | 4 -- mlir/lib/Dialect/GPU/TransformOps/CMakeLists.txt | 1 - .../Dialect/GPU/TransformOps/GPUTransformOps.cpp | 2 - .../lib/Dialect/MemRef/TransformOps/CMakeLists.txt | 1 - .../MemRef/TransformOps/MemRefTransformOps.cpp | 2 - .../lib/Dialect/Tensor/TransformOps/CMakeLists.txt | 1 - .../transform-op-simplify-bounded-affine-ops.mlir | 30 ++++----- .../Bufferization/Transforms/transform-ops.mlir | 38 +++++------ mlir/test/Dialect/GPU/transform-gpu-failing.mlir | 58 ++++++++--------- mlir/test/Dialect/GPU/transform-gpu.mlir | 44 ++++++------- .../MemRef/extract-address-computations.mlir | 74 +++++++++++----------- .../test/Dialect/MemRef/make-loop-independent.mlir | 12 ++-- mlir/test/Dialect/MemRef/transform-ops.mlir | 52 +++++++-------- .../Tensor/transform-op-make-loop-independent.mlir | 30 ++++----- utils/bazel/llvm-project-overlay/mlir/BUILD.bazel | 9 --- 27 files changed, 193 insertions(+), 221 deletions(-) diff --git a/mlir/include/mlir/Dialect/Affine/TransformOps/AffineTransformOps.h b/mlir/include/mlir/Dialect/Affine/TransformOps/AffineTransformOps.h index 4a27f30..fb6f96e 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 76f0d8e..b74e4af 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:$bounded_values, + let arguments = (ins TransformHandleTypeInterface:$target, + Variadic:$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 06204b6..426e9d0 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 692f0f0..49b5ef6 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 57d74d8..3b10fcb 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 c719fed..4216231 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:$block_dims, DefaultValuedOptionalAttr:$warp_dims, DefaultValuedAttr:$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:$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 dd33df9..3d85a1b 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 1cefe56..04dfe1f 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:$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 c735700..e902cce 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 42be882..a9580bd 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 9a952e6..5a7f092 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 24e2c837..4797cd9 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 6dd32b8..9d45442 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(); - declareGeneratedDialect(); declareGeneratedDialect(); diff --git a/mlir/lib/Dialect/GPU/TransformOps/CMakeLists.txt b/mlir/lib/Dialect/GPU/TransformOps/CMakeLists.txt index 61695b1..7abd5b9 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 8cd2ccf..e9027d1 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(); declareGeneratedDialect(); declareGeneratedDialect(); declareGeneratedDialect(); diff --git a/mlir/lib/Dialect/MemRef/TransformOps/CMakeLists.txt b/mlir/lib/Dialect/MemRef/TransformOps/CMakeLists.txt index b32e06a..0c9ee3d 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 c523af9..f8b4491 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(); declareGeneratedDialect(); declareGeneratedDialect(); declareGeneratedDialect(); diff --git a/mlir/lib/Dialect/Tensor/TransformOps/CMakeLists.txt b/mlir/lib/Dialect/Tensor/TransformOps/CMakeLists.txt index be1a5dd..27ea475 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 2f384ea..a30607b 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 05d6a1b..c4a4044 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, %v : vector<4xf32>) -> (tensor !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, %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) { @@ -69,9 +69,9 @@ func.func @test_unknown_op_failure() -> (tensor) { // ----- 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 813b6f3..0d56048 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, #gpu.thread, #gpu.thread ] ) : (!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 d9872a9..64d80f9 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 17e2ac3..5064f60 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>) } 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, } 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, } 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, } 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, } 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 !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, } 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 86ce9f4..1a34d9c 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 3edcc53..8d2d569 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 18a99c5..d379888 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 } diff --git a/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel b/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel index 603716c..1c06433 100644 --- a/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel +++ b/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel @@ -1210,7 +1210,6 @@ cc_library( ":AffineUtils", ":FuncDialect", ":IR", - ":PDLDialect", ":TransformDialect", ":Transforms", ":VectorDialect", @@ -4385,7 +4384,6 @@ td_library( ], includes = ["include"], deps = [ - ":PDLDialectTdFiles", ":TransformDialectTdFiles", ], ) @@ -4430,7 +4428,6 @@ cc_library( ":GPUTransformOpsIncGen", ":GPUTransforms", ":IR", - ":PDLDialect", ":Parser", ":SCFDialect", ":SideEffectInterfaces", @@ -5920,7 +5917,6 @@ td_library( ], includes = ["include"], deps = [ - ":PDLDialect", ":TransformDialectTdFiles", ], ) @@ -5953,7 +5949,6 @@ cc_library( deps = [ ":AffineDialect", ":IR", - ":PDLDialect", ":SCFDialect", ":TensorDialect", ":TensorTransformOpsIncGen", @@ -10682,7 +10677,6 @@ td_library( ], includes = ["include"], deps = [ - ":PDLDialect", ":TransformDialectTdFiles", ], ) @@ -10721,7 +10715,6 @@ cc_library( ":MemRefTransformOpsIncGen", ":MemRefTransforms", ":NVGPUDialect", - ":PDLDialect", ":SCFDialect", ":TransformDialect", ":TransformUtils", @@ -10956,7 +10949,6 @@ td_library( includes = ["include"], deps = [ ":BufferizationOpsTdFiles", - ":PDLDialectTdFiles", ":TransformDialectTdFiles", ], ) @@ -10997,7 +10989,6 @@ cc_library( ":BufferizationTransforms", ":IR", ":MemRefDialect", - ":PDLDialect", ":Parser", ":SideEffectInterfaces", ":TensorDialect", -- 2.7.4