[mlir] don't hardcode PDL_Operation in Transform dialect extensions

[ftynse]

Summary

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.

[ftynse]

Also do Bufferization dialect

mlir/include/mlir/Dialect/Affine/TransformOps/AffineTransformOps.h

 //===- AffineTransformOps.h - Affine transformation ops ---------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #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"
 
 namespace mlir {
 namespace func {
 class FuncOp;
 } // namespace func

mlir/include/mlir/Dialect/Affine/TransformOps/AffineTransformOps.td

 //===- AffineTransformOps.td - Affine transformation ops ---*- tablegen -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #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"
 include "mlir/Interfaces/SideEffectInterfaces.td"
 include "mlir/IR/OpBase.td"
 
 def Transform_AffineForOp : Transform_ConcreteOpType<"affine.for">;
 
let description = [{
     Target ops must be affine.min or affine.max ops. This transform consumes the
     target handle and does not produce any handle. It reads the bounded op
     handles.
 
     TODO: Support affine.apply targets.
     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)
   }];
 }
 
 #endif // Affine_TRANSFORM_OPS

mlir/include/mlir/Dialect/Bufferization/TransformOps/BufferizationTransformOps.h

 //===- BufferizationTransformOps.h - Buff. transf. ops ----------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef MLIR_DIALECT_BUFFERIZATION_TRANSFORMOPS_BUFFERIZATIONTRANSFORMOPS_H
 #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"
 
 namespace mlir {
 namespace tensor {
 class EmptyOp;
 } // namespace tensor

mlir/include/mlir/Dialect/Bufferization/TransformOps/BufferizationTransformOps.td

 //===- BufferizationTransformOps.td - Buff. transf. ops ----*- tablegen -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef BUFFERIZATION_TRANSFORM_OPS
 #define BUFFERIZATION_TRANSFORM_OPS
 
 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"
 
 def Transform_EmptyOp : Transform_ConcreteOpType<"tensor.empty">;
 def Transform_AllocTensorOp : Transform_ConcreteOpType<"bufferization.alloc_tensor">;
 
 //===----------------------------------------------------------------------===//
 // OneShotBufferizeOp
let description = [{
     See `-eliminate-empty-tensors` for more details.
 
     #### Return modes
 
     This transform reads the target handle and modifies the payload. It does
     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)";
 }
 
 //===----------------------------------------------------------------------===//
 // EmptyTensorToAllocTensorOp
 //===----------------------------------------------------------------------===//
 
 def EmptyTensorToAllocTensorOp
     : Op<Transform_Dialect, "bufferization.empty_tensor_to_alloc_tensor",

mlir/include/mlir/Dialect/GPU/TransformOps/GPUTransformOps.h

 //===- GPUTransformOps.h - GPU transform ops --------------------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #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"
 #include "mlir/IR/PatternMatch.h"
 
 namespace mlir {
 namespace gpu {
 class GpuOp;

mlir/include/mlir/Dialect/GPU/TransformOps/GPUTransformOps.td

 //===- GPUTransformOps.td - GPU transform ops --------------*- tablegen -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef GPU_TRANSFORM_OPS
 #define GPU_TRANSFORM_OPS
 
 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"
 
 def MapNestedForallToThreads :
   Op<Transform_Dialect, "gpu.map_nested_forall_to_threads",
     [FunctionalStyleTransformOpTrait,
      MemoryEffectsOpInterface,
      TransformEachOpTrait,
let description = [{
           threads in each dimension. This is a mandatory attribute that is used
           to constrain the number of threads in each dimension. If an 
           `scf.forall` op is mapped to fewer threads, predication occurs.
         - the warp_dims is a list of integers that specifies the number of
           warps in each dimension. This is an optional attribute that is used
           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.
 
       Only **bufferized** `scf.forall` are currently supported.
       Only `scf.forall` distributed to **at most 3 dimensions** are
       currently supported.
 
let description = [{
           ... // body 2
         }
         gpu.barrier
         gpu.terminator
       }
       ```
     }];
 
-  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
     `block_dims` `=` $block_dims
     (`warp_dims` `=` $warp_dims^)?
     (`sync_after_distribute` `=` $sync_after_distribute^)?
     attr-dict
+    `:` functional-type($target, $result)
   }];
   let extraClassDeclaration = [{
     ::mlir::DiagnosedSilenceableFailure applyToOne(
         ::mlir::Operation *target,
         ::mlir::transform::ApplyToEachResultList &results,
         ::mlir::transform::TransformState &state);
   }];
 }
let description = [{
     referred to by the `target` PDLOperation lower to GPU properly, the
     transform succeeds. Otherwise the transform definitely fails.
 
     The returned handle points to the same LaunchOp operand, consuming it and
     producing a new SSA value to satisfy chaining and linearity of the IR
     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(
         ::mlir::Operation *target,
         ::mlir::transform::ApplyToEachResultList &results,
         ::mlir::transform::TransformState &state);
   }];
 }
 
 #endif // GPU_TRANSFORM_OPS

mlir/include/mlir/Dialect/MemRef/TransformOps/MemRefTransformOps.h

 //===- MemRefTransformOps.h - MemRef transformation ops ---------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #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"
 
 namespace mlir {
 namespace memref {
 class AllocOp;
 } // namespace memref
 } // namespace mlir

mlir/include/mlir/Dialect/MemRef/TransformOps/MemRefTransformOps.td

 //===- MemRefTransformOps.td - MemRef transformation ops --*- tablegen -*--===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef MEMREF_TRANSFORM_OPS
 #define MEMREF_TRANSFORM_OPS
 
 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"
 
 def Transform_MemRefAllocOp : Transform_ConcreteOpType<"memref.alloc">;
 
 def MemRefMultiBufferOp : Op<Transform_Dialect, "memref.multibuffer",
     [FunctionalStyleTransformOpTrait, MemoryEffectsOpInterface,
      DeclareOpInterfaceMethods<TransformOpInterface>]> {
let description = [{
      succeeds, and failure otherwise.
   }];
 
   let arguments =
       (ins Transform_MemRefAllocOp:$target,
            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)";
 }
 
 def MemRefExtractAddressComputationsOp :
     Op<Transform_Dialect, "memref.extract_address_computations",
        [FunctionalStyleTransformOpTrait, MemoryEffectsOpInterface,
let description = [{
     #### Return modes
 
     This operation produces `definiteFailure` if the extraction fails for any
     reason.
     The operation always returns the handle to the target op that is expected
     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)";
 
   let extraClassDeclaration = [{
     ::mlir::DiagnosedSilenceableFailure applyToOne(
         ::mlir::Operation *target,
         ::mlir::transform::ApplyToEachResultList &transformResults,
         ::mlir::transform::TransformState &state);
let description = [{
     variables, this transform succeeds and returns the unmodified target op.
 
     Otherwise, the returned handle points to a subset of the produced ops:
     - memref.alloca: The returned handle points to the memref.subview op.
 
     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(
         ::mlir::Operation *target,
         ::mlir::transform::ApplyToEachResultList &results,
         ::mlir::transform::TransformState &state);
   }];
 }
 
 #endif // MEMREF_TRANSFORM_OPS

mlir/include/mlir/Dialect/Tensor/TransformOps/TensorTransformOps.h

 //===- TensorTransformOps.h - Tensor transformation ops ---------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #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"
 #include "mlir/IR/PatternMatch.h"
 
 namespace mlir {
 class DialectRegistry;
 

mlir/include/mlir/Dialect/Tensor/TransformOps/TensorTransformOps.td

 //===- TensorTransformOps.td - Tensor transformation ops ---*- tablegen -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #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"
 include "mlir/Interfaces/SideEffectInterfaces.td"
 include "mlir/IR/OpBase.td"
 
 def Transform_TensorPadOp : Transform_ConcreteOpType<"tensor.pad">;
 
let description = [{
 
     Otherwise, the returned handle points to a subset of the produced ops:
     - tensor.empty: The returned handle points to the tensor.extract_slice op.
     - tensor.pad: The returned handle points to the tensor.extract_slice op.
 
     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(
         ::mlir::Operation *target,
         ::mlir::transform::ApplyToEachResultList &results,
         ::mlir::transform::TransformState &state);
   }];
 }
 
 #endif // TENSOR_TRANSFORM_OPS

mlir/lib/Dialect/Affine/TransformOps/AffineTransformOps.cpp

 //=== AffineTransformOps.cpp - Implementation of Affine transformation ops ===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/Affine/TransformOps/AffineTransformOps.h"
 #include "mlir/Dialect/Affine/Analysis/AffineStructures.h"
 #include "mlir/Dialect/Affine/Analysis/Utils.h"
 #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"
 
 using namespace mlir;
 using namespace mlir::affine;
 using namespace mlir::transform;
 

mlir/lib/Dialect/Affine/TransformOps/CMakeLists.txt

 add_mlir_dialect_library(MLIRAffineTransformOps
   AffineTransformOps.cpp
 
   ADDITIONAL_HEADER_DIRS
   ${MLIR_MAIN_INCLUDE_DIR}/mlir/Dialect/Affine/TransformOps
 
   DEPENDS
   MLIRAffineTransformOpsIncGen
 
   LINK_LIBS PUBLIC
   MLIRAffineAnalysis
   MLIRAffineDialect
   MLIRFuncDialect
   MLIRIR
-  MLIRPDLDialect
   MLIRAffineDialect
   MLIRAffineTransforms
   MLIRAffineUtils
   MLIRTransformDialect
   MLIRVectorDialect
 )

mlir/lib/Dialect/Bufferization/TransformOps/BufferizationTransformOps.cpp

 //===- BufferizationTransformOps.h - Bufferization transform ops ----------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/Bufferization/TransformOps/BufferizationTransformOps.h"
 
 #include "mlir/Dialect/Bufferization/IR/Bufferization.h"
 #include "mlir/Dialect/Bufferization/Transforms/OneShotAnalysis.h"
 #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"
 
 using namespace mlir;
 using namespace mlir::bufferization;
 using namespace mlir::transform;
 

 /// ops are using PDL types for operands and results.
 class BufferizationTransformDialectExtension
     : public transform::TransformDialectExtension<
           BufferizationTransformDialectExtension> {
 public:
   using Base::Base;
 
   void init() {
-    declareDependentDialect<pdl::PDLDialect>();
-
     declareGeneratedDialect<bufferization::BufferizationDialect>();
     declareGeneratedDialect<memref::MemRefDialect>();
 
     registerTransformOps<
 #define GET_OP_LIST
 #include "mlir/Dialect/Bufferization/TransformOps/BufferizationTransformOps.cpp.inc"
         >();
   }

mlir/lib/Dialect/GPU/TransformOps/CMakeLists.txt

 add_mlir_dialect_library(MLIRGPUTransformOps
   GPUTransformOps.cpp
 
   ADDITIONAL_HEADER_DIRS
   ${MLIR_MAIN_INCLUDE_DIR}/mlir/Dialect/GPU/TransformOps
   ${MLIR_MAIN_INCLUDE_DIR}/mlir/Interfaces
 
   DEPENDS
   MLIRGPUTransformOpsIncGen
   MLIRDeviceMappingInterfacesIncGen
   MLIRGPUDeviceMapperEnumsGen
   
   LINK_LIBS PUBLIC
   MLIRIR
   MLIRGPUTransforms
   MLIRParser
-  MLIRPDLDialect
   MLIRSideEffectInterfaces
   MLIRTransformDialect
   MLIRGPUDialect
   )  

mlir/lib/Dialect/GPU/TransformOps/GPUTransformOps.cpp

 //===- GPUTransformOps.cpp - Implementation of GPU transform ops ----------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/GPU/TransformOps/GPUTransformOps.h"
 
 #include "mlir/Dialect/Affine/IR/AffineOps.h"
 #include "mlir/Dialect/Arith/IR/Arith.h"
 #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"
 #include "mlir/Dialect/Transform/IR/TransformInterfaces.h"
 #include "mlir/Dialect/Utils/IndexingUtils.h"
 #include "mlir/IR/AffineExpr.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/BuiltinAttributes.h"

 namespace {
 /// Registers new ops and declares PDL as dependent dialect since the
 /// additional ops are using PDL types for operands and results.
 class GPUTransformDialectExtension
     : public transform::TransformDialectExtension<
           GPUTransformDialectExtension> {
 public:
   GPUTransformDialectExtension() {
-    declareDependentDialect<pdl::PDLDialect>();
     declareGeneratedDialect<scf::SCFDialect>();
     declareGeneratedDialect<arith::ArithDialect>();
     declareGeneratedDialect<GPUDialect>();
     registerTransformOps<
 #define GET_OP_LIST
 #include "mlir/Dialect/GPU/TransformOps/GPUTransformOps.cpp.inc"
         >();
   }

mlir/lib/Dialect/MemRef/TransformOps/CMakeLists.txt

 add_mlir_dialect_library(MLIRMemRefTransformOps
   MemRefTransformOps.cpp
 
   ADDITIONAL_HEADER_DIRS
   ${MLIR_MAIN_INCLUDE_DIR}/mlir/Dialect/MemRef/TransformOps
 
   DEPENDS
   MLIRMemRefTransformOpsIncGen
 
   LINK_LIBS PUBLIC
   MLIRAffineDialect
   MLIRArithDialect
   MLIRIR
-  MLIRPDLDialect
   MLIRLoopLikeInterface
   MLIRMemRefDialect
   MLIRMemRefTransforms
   MLIRNVGPUDialect
   MLIRTransformDialect
   MLIRVectorDialect
 )

mlir/lib/Dialect/MemRef/TransformOps/MemRefTransformOps.cpp

 //===- MemRefTransformOps.cpp - Implementation of Memref transform ops ----===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/MemRef/TransformOps/MemRefTransformOps.h"
 #include "mlir/Dialect/Affine/IR/AffineOps.h"
 #include "mlir/Dialect/Arith/IR/Arith.h"
 #include "mlir/Dialect/MemRef/IR/MemRef.h"
 #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"
 #include "mlir/Dialect/Vector/IR/VectorOps.h"
 #include "mlir/Interfaces/LoopLikeInterface.h"
 #include "mlir/Transforms/GreedyPatternRewriteDriver.h"
 #include "llvm/Support/Debug.h"
 

 namespace {
 class MemRefTransformDialectExtension
     : public transform::TransformDialectExtension<
           MemRefTransformDialectExtension> {
 public:
   using Base::Base;
 
   void init() {
-    declareDependentDialect<pdl::PDLDialect>();
     declareGeneratedDialect<affine::AffineDialect>();
     declareGeneratedDialect<arith::ArithDialect>();
     declareGeneratedDialect<memref::MemRefDialect>();
     declareGeneratedDialect<nvgpu::NVGPUDialect>();
     declareGeneratedDialect<vector::VectorDialect>();
 
     registerTransformOps<
 #define GET_OP_LIST

mlir/lib/Dialect/Tensor/TransformOps/CMakeLists.txt

 add_mlir_dialect_library(MLIRTensorTransformOps
   TensorTransformOps.cpp
 
   ADDITIONAL_HEADER_DIRS
   ${MLIR_MAIN_INCLUDE_DIR}/mlir/Dialect/Tensor/TransformOps
 
   DEPENDS
   MLIRTensorTransformOpsIncGen
 
   LINK_LIBS PUBLIC
   MLIRAffineDialect
   MLIRIR
-  MLIRPDLDialect
   MLIRSCFDialect
   MLIRTensorDialect
   MLIRTensorTransforms
   MLIRTransformDialect
 )

mlir/test/Dialect/Affine/transform-op-simplify-bounded-affine-ops.mlir

 // RUN: mlir-opt  %s -allow-unregistered-dialect \
 // RUN:     --test-transform-dialect-interpreter -verify-diagnostics \
 // RUN:     --split-input-file | FileCheck %s
 
 //     CHECK: func @simplify_min_max()
 // CHECK-DAG:   %[[c50:.*]] = arith.constant 50 : index
 // CHECK-DAG:   %[[c100:.*]] = arith.constant 100 : index
 //     CHECK:   return %[[c50]], %[[c100]]
 func.func @simplify_min_max() -> (index, index) {
   %0 = "test.some_op"() : () -> (index)
   %1 = affine.min affine_map<()[s0] -> (50, 100 - s0)>()[%0]
   %2 = affine.max affine_map<()[s0] -> (100, 80 + s0)>()[%0]
   return %1, %2 : 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
 }
 
 // -----
 
 // CHECK: func @simplify_min_sequence()
 // CHECK:   %[[c1:.*]] = arith.constant 1 : index
 // CHECK:   return %[[c1]]
 func.func @simplify_min_sequence() -> index {
   %1 = "test.workgroup_id"() : () -> (index)
   %2 = affine.min affine_map<()[s0] -> (s0 * -32 + 1023, 32)>()[%1]
   %3 = "test.thread_id"() : () -> (index)
   %4 = affine.min affine_map<()[s0, s1] -> (s0 - s1 * (s0 ceildiv 32), s0 ceildiv 32)>()[%2, %3]
   return %4 : 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
 }

mlir/test/Dialect/Bufferization/Transforms/transform-ops.mlir

 // RUN: mlir-opt --test-transform-dialect-interpreter %s -split-input-file -verify-diagnostics | FileCheck %s
 
 // 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(
 //  CHECK-SAME:     %[[A:.*]]: tensor<?xf32>
 func.func @test_function(%A : tensor<?xf32>, %v : vector<4xf32>) -> (tensor<?xf32>) {
   %c0 = arith.constant 0 : index
 
   // CHECK: %[[A_memref:.*]] = bufferization.to_memref %[[A]]
func.func @test_function(%A : tensor<?xf32>, %v : vector<4xf32>) -> (tensor<?xf32>) {
   return %0 : tensor<?xf32>
 }
 
 // -----
 
 // 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(
 //  CHECK-SAME:     %[[A:.*]]: tensor<?xf32>
 func.func @test_function_analysis(%A : tensor<?xf32>, %v : vector<4xf32>) -> (tensor<?xf32>) {
   %c0 = arith.constant 0 : index
   // CHECK: vector.transfer_write
   // CHECK-SAME: {__inplace_operands_attr__ = ["none", "false", "none"]}
   // CHECK-SAME: tensor<?xf32>
   %0 = vector.transfer_write %v, %A[%c0] : vector<4xf32>, tensor<?xf32>
   return %0 : tensor<?xf32>
 }
 
 // -----
 
 // Test One-Shot Bufferize transform failure with an unknown op. This would be
 // 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>) {
   // expected-error @+1 {{op was not bufferized}}
   %0 = "test.dummy_op"() : () -> (tensor<?xf32>)
   return %0 : 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 {
   // CHECK-LABEL: func @test_function(
   //  CHECK-SAME:     %[[A:.*]]: tensor<?xf32>
   func.func @test_function(%A : tensor<?xf32>, %v : vector<4xf32>) -> (tensor<?xf32>) {
     %c0 = arith.constant 0 : index
 
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(
 // CHECK-SAME:  %[[A:.*]]: memref<12x9xf32>,
 // CHECK-SAME:  %[[B:.*]]: memref<9x6xf32>,
 // CHECK-SAME:  %[[C:.*]]: memref<12x6xf32>) -> memref<12x6xf32> {
 func.func @matmul(%A: tensor<12x9xf32>, %B: tensor<9x6xf32>, %C: tensor<12x6xf32>) -> tensor<12x6xf32> {
   // CHECK: linalg.matmul ins(%[[A]], %[[B]] : memref<12x9xf32>, memref<9x6xf32>) outs(%[[C]] : memref<12x6xf32>)
   %D = linalg.matmul ins(%A, %B: tensor<12x9xf32>, tensor<9x6xf32>) outs(%C: tensor<12x6xf32>) -> tensor<12x6xf32>
   // CHECK: return %[[C]] : memref<12x6xf32>
   return %D : 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">
 }
 
 // Expect `bufferization.empty_tensor_to_alloc_tensor` to replace the tensor.empty.
 func.func @empty_to_tensor_alloc() -> tensor<2x2xf32> {
   // CHECK: bufferization.alloc_tensor
   %0 = tensor.empty() : tensor<2x2xf32>
   return %0 : 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(
 //       CHECK:   tensor.extract_slice
 //       CHECK:   linalg.fill
 //       CHECK:   tensor.insert_slice
 func.func @empty_tensor_elimination(
     %t: tensor<10xf32>, %f: f32) -> tensor<10xf32> {
   %0 = tensor.empty() : tensor<5xf32>
   %1 = linalg.fill ins(%f : f32) outs(%0 : tensor<5xf32>) -> tensor<5xf32>
   %2 = tensor.insert_slice %1 into %t [1][5][1]
       : tensor<5xf32> into tensor<10xf32>
   return %2 : tensor<10xf32>
 }

mlir/test/Dialect/GPU/transform-gpu-failing.mlir

 // RUN: mlir-opt --test-transform-dialect-interpreter --split-input-file  -canonicalize -cse --verify-diagnostics %s
 
 func.func @map_nested_forall_to_threads_not_gpu_launch() -> () {
   %1 = tensor.empty() : tensor<4xf32>
   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
 }
 
 // -----
 
 func.func @map_nested_forall_to_threads_excessive_threads(%x: memref<2 x 32 x f32>, %y: memref<2 x 32 x f32>, %t: memref<32 x f32>, %alpha : f32, %stream : !gpu.async.token) -> memref<2 x 32 x f32> {
   %one = arith.constant 1 : index
   %c900 = arith.constant 900 : index
   %c9 = arith.constant 9 : index
scf.forall (%i, %j) in (%c7, %c9) {
         memref.store %6, %y[%i, %j] : memref<2 x 32 x f32>
      }  { mapping = [#gpu.thread<y>, #gpu.thread<x>] }
     gpu.terminator
   }
 
   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
 }
 
 // -----
 
 func.func @map_nested_forall_to_threads_fewer_threads(%x: memref<2 x 32 x f32>, %y: memref<2 x 32 x f32>, %t: memref<32 x f32>, %alpha : f32, %stream : !gpu.async.token) -> memref<2 x 32 x f32> {
   %one = arith.constant 1 : index
   %c900 = arith.constant 900 : index
   %c9 = arith.constant 9 : index
%name2 = gpu.launch async[%stream] blocks(%arg3, %arg4, %arg5) in (%arg9 = %one, %arg10 = %one, %arg11 = %one)
      }  { mapping = [#gpu.thread<y>, #gpu.thread<x>] }
     gpu.terminator
   }
 
   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 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
 }
 
 // -----
 
 func.func @map_nested_forall_to_threads_dynamic_trip_count(%x: memref<2 x 32 x f32>, %y: memref<2 x 32 x f32>, %t: memref<32 x f32>, %alpha : f32, %stream : !gpu.async.token, %c9 : index, %c7 : index) -> memref<2 x 32 x f32> {
   %one = arith.constant 1 : index
   %c900 = arith.constant 900 : index
   %name = gpu.launch async[%stream] blocks(%arg3, %arg4, %arg5) in (%arg9 = %one, %arg10 = %one, %arg11 = %one)
             threads(%arg6, %arg7, %arg8) in (%arg12 = %one, %arg13 = %one, %arg14 = %one)
   {
     scf.forall (%i, %j) in (%c7, %c900) {
         %4 = memref.load %x[%i, %j] : memref<2 x 32 x f32>
         %5 = memref.load %y[%i, %j] : memref<2 x 32 x f32>
         %6 = math.fma %alpha, %4, %5 : f32
         memref.store %6, %y[%i, %j] : memref<2 x 32 x f32>
      }  { mapping = [#gpu.thread<y>, #gpu.thread<x>] }
     gpu.terminator
   }
   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 {{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
 }
 
 // -----
 
 func.func @map_nested_forall_to_threads_not_buffer(%x: tensor<32x32xf32>, %y: tensor<32x32xf32>, %z: tensor<32x32xf32>, %stream : !gpu.async.token) {
   %one = arith.constant 1 : index
   %name = gpu.launch async[%stream] blocks(%arg3, %arg4, %arg5) in (%arg9 = %one, %arg10 = %one, %arg11 = %one)
             threads(%arg6, %arg7, %arg8) in (%arg12 = %one, %arg13 = %one, %arg14 = %one)
   {
     %t = linalg.matmul ins(%x, %y: tensor<32x32xf32>, tensor<32x32xf32>) outs(%z : tensor<32x32xf32>) -> tensor<32x32xf32>
     gpu.terminator
   }
   return
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%arg0: !transform.any_op):
   %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
 }
 
 // -----
 
 
 func.func @map_forall_to_blocks_not_gpu_launch() -> () {
   // expected-note @below {{when applied to this payload op}}
   %1 = tensor.empty() : tensor<4xf32>
   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
 }
 
 // -----
 
 func.func @map_forall_to_blocks_not_unique(%x: memref<2 x 32 x f32>, %y: memref<2 x 32 x f32>, %t: memref<32 x f32>, %alpha : f32, %stream : !gpu.async.token) -> memref<2 x 32 x f32> {
   %one = arith.constant 1 : index
   %c900 = arith.constant 900 : index
   %c9 = arith.constant 9 : index
%name = gpu.launch async[%stream] blocks(%arg3, %arg4, %arg5) in (%arg9 = %one, %arg10 = %one, %arg11 = %one)
      }  { mapping = [#gpu.thread<y>, #gpu.thread<x>] }
     gpu.terminator
   }
 
   return %y : memref<2 x 32 x f32>
 }
 
 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
 }
 
 // -----
 
 // expected-note @below {{when applied to this payload op}}
 func.func @map_forall_to_blocks_large_loop(%x: memref<2 x 32 x f32>, %y: memref<2 x 32 x f32>, %t: memref<32 x f32>, %alpha : f32, %stream : !gpu.async.token) -> memref<2 x 32 x f32> {
   %one = arith.constant 1 : index
   %c65537 = arith.constant 65536 : index
scf.forall (%i, %j) in (%c7, %c9) {
       %6 = math.fma %alpha, %4, %5 : f32
       memref.store %6, %y[%i, %j] : memref<2 x 32 x f32>
   }  { mapping = [#gpu.thread<y>, #gpu.thread<x>] }
 
   return %y : memref<2 x 32 x f32>
 }
 
 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
 }
 
 // -----
 
 func.func @map_forall_to_blocks_large_loop(%x: memref<2 x 32 x f32>, %y: memref<2 x 32 x f32>, %t: memref<32 x f32>, %alpha : f32, %stream : !gpu.async.token) -> memref<2 x 32 x f32> {
   %one = arith.constant 1 : index
   %c65535 = arith.constant 65535 : index
   scf.forall (%i, %j) in (%c65535, %c65535) {
       %4 = memref.load %x[%i, %j] : memref<2 x 32 x f32>
       %5 = memref.load %y[%i, %j] : memref<2 x 32 x f32>
       %6 = math.fma %alpha, %4, %5 : f32
       memref.store %6, %y[%i, %j] : memref<2 x 32 x f32>
   }  { mapping = [#gpu.block<x>, #gpu.block<y>] }
   return %y : memref<2 x 32 x f32>
 }
 
 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
 }
 
 // -----
 
 !type = memref<32x32xf32>
 func.func @saxpy2d_singleloop(%x: !type, %y: !type, %stream : !gpu.async.token) -> !type {
   %c32 = arith.constant 32 : index
   %one = arith.constant 1 : index
   %name = gpu.launch async[%stream] blocks(%arg3, %arg4, %arg5) in (%arg9 = %one, %arg10 = %one, %arg11 = %one)
             threads(%arg6, %arg7, %arg8) in (%arg12 = %one, %arg13 = %one, %arg14 = %one)
   {
     scf.forall (%i, %j) in (%c32, %c32) {
         %4 = memref.load %x[%i, %j] : !type
         %5 = memref.load %y[%i, %j] : !type
         %6 = arith.mulf %4, %5 : f32
         memref.store %6, %y[%i, %j] : !type
      }  { mapping = [#gpu.thread<x>, #gpu.thread<x>] }
     gpu.terminator
   }
   return %y : !type
 }
 
 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
 }
 
 // -----
 
 func.func @tiling_buffer_semantic_op(%x: memref<32x32xf32>, %y: memref<32x32xf32>, %stream : !gpu.async.token) {
   %one = arith.constant 1 : index
   %name = gpu.launch async[%stream] blocks(%arg3, %arg4, %arg5) in (%arg9 = %one, %arg10 = %one, %arg11 = %one)
             threads(%arg6, %arg7, %arg8) in (%arg12 = %one, %arg13 = %one, %arg14 = %one)

mlir/test/Dialect/GPU/transform-gpu.mlir

scf.forall (%i, %j) in (%c7, %c9) {
         memref.store %6, %y[%i, %j] : !type
      }  { mapping = [#gpu.block<x>, #gpu.block<y>]}
     gpu.terminator
   }
   return %y : !type
 }
 
 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
 }
 
 // -----
 
 !type = memref<2 x 32 x f32>
 !type1d = memref<32 x f32>
 
 // CHECK-LABEL: func.func @saxpy2d(
scf.forall (%i) in (%c12) {
         memref.store %8, %t[%i] : !type1d
      }  {mapping = [#gpu.thread<x>] }
     gpu.terminator
   }
   return %y : !type
 }
 
 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
 }
 
 // -----
 
 !type4d = memref<32x64x4x32xf32>
 
 // CHECK-LABEL: func.func @saxpy4d(
 // CHECK-SAME:    %[[ARGX:[0-9a-z]+]]: memref<32x64x4x32xf32>
scf.forall (%k, %l) in (%c4, %c32) {
       %6 = math.fma %alpha, %4, %5 : f32
       memref.store %6, %y[%i, %j, %k, %l] : !type4d
     }  { mapping = [#gpu.thread<y>, #gpu.thread<x>] }
   }  { mapping = [#gpu.block<x>, #gpu.block<y>] }
   return %y : !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
 }
 
 // -----
 
 !type = memref<2 x 32 x f32>
 !type1d = memref<32 x f32>
 
 // CHECK-LABEL: func.func @saxpy2d_no_barrier(
scf.forall (%i, %j) in (%c7, %c9) {
         memref.store %6, %y[%i, %j] : !type
      }  { mapping = [#gpu.thread<y>, #gpu.thread<x>] }
     gpu.terminator
   }
   return %y : !type
 }
 
 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
 }
 
 // -----
 
 !type = memref<32x32xf32>
 // CHECK-LABEL: func.func @saxpy2d_singleloop(
 // CHECK-SAME:    %[[ARGX:[0-9a-z]+]]: memref<32x32xf32>
 // CHECK-SAME:    %[[ARGY:[0-9a-z]+]]: memref<32x32xf32>
scf.forall (%i) in (%c32) {
         memref.store %6, %y[%i, %i] : !type
      }  { mapping = [#gpu.thread<x>] }
     gpu.terminator
   }
   return %y : !type
 }
 
 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
 }
 
 // -----
 
 !type = memref<3 x 2 x 32 x f32>
 !type1d = memref<32 x f32>
 
 // CHECK-LABEL: func.func @saxpy3d_fold_id_z(
//      CHECK:   memref.store %{{.*}}, %{{.*}}[%[[C0]]
         memref.store %6, %y[%i, %j, %k] : !type
      }  { mapping = [#gpu.thread<z>, #gpu.thread<y>, #gpu.thread<x>] }
     gpu.terminator
   }
   return %y : !type
 }
 
 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
 }
 
 // -----
 
 !type = memref<2 x 32 x f32>
 !type1d = memref<32 x f32>
 
 // CHECK-DAG: #[[$MAPWX:.*]] = affine_map<(d0, d1) -> (((d0 + d1 * 12) floordiv 32) mod 3)>
scf.forall (%i, %j) in (%c10, %c2) {
         memref.store %8, %t[%j] : !type1d
      }  {mapping = [#gpu.linear<x>, #gpu.linear<y>] }
     gpu.terminator
   }
   return %y : !type
 }
 
 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
 }

mlir/test/Dialect/MemRef/extract-address-computations.mlir

 func.func @test_load(%base : memref<2x16x16xf32>, %offset : index) -> f32 {
   %c0 = arith.constant 0 : index
   %c8 = arith.constant 8 : index
   %loaded_val = memref.load %base[%offset, %c0, %c8] : memref<2x16x16xf32>
   return %loaded_val : 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
 }
 
 // -----
 
 // Same as previous @test_load but with the nontemporal flag.
 
 // CHECK-LABEL: @test_load_nontemporal(
 // CHECK-SAME: %[[BASE:[^:]*]]: memref{{[^,]*}},
 // CHECK-SAME: %[[DYN_OFFSET:.*]]: index)
 // CHECK-DAG: %[[C0:.*]] = arith.constant 0 : index
 // CHECK-DAG: %[[SUBVIEW:.*]] = memref.subview %[[BASE]][%[[DYN_OFFSET]], 0, 8] [1, 1, 1] [1, 1, 1] : memref<2x16x16xf32> to memref<1x1x1xf32, strided<[256, 16, 1], offset: ?>>
 // CHECK: %[[LOADED_VAL:.*]] = memref.load %[[SUBVIEW]][%[[C0]], %[[C0]], %[[C0]]] {nontemporal = true} : memref<1x1x1xf32, strided<[256, 16, 1], offset: ?>>
 // CHECK: return %[[LOADED_VAL]] : f32
 func.func @test_load_nontemporal(%base : memref<2x16x16xf32>, %offset : index) -> f32 {
   %c0 = arith.constant 0 : index
   %c8 = arith.constant 8 : index
   %loaded_val = memref.load %base[%offset, %c0, %c8] {nontemporal = true } : memref<2x16x16xf32>
   return %loaded_val : 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
 }
 
 // -----
 
 // Simple test: check that we extract the address computation of a store into
 // a dedicated subview.
 // The resulting store will use the address from the subview and have only
 // indices set to zero.
func.func @test_store(%base : memref<2x16x16xf32>, %offset : index) -> () {
   %cf0 = arith.constant 0.0 : f32
   %c0 = arith.constant 0 : index
   %c8 = arith.constant 8 : index
   memref.store %cf0, %base[%offset, %c0, %c8] : memref<2x16x16xf32>
   return
 }
 
 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
 }
 
 // -----
 
 // Same as @test_store but check that the nontemporal flag is preserved.
 
 // CHECK-LABEL: @test_store_nontemporal(
 // CHECK-SAME: %[[BASE:[^:]*]]: memref{{[^,]*}},
 // CHECK-SAME: %[[DYN_OFFSET:.*]]: index)
 // CHECK-DAG: %[[CF0:.*]] = arith.constant 0.0{{0*e\+00}} : f32
 // CHECK-DAG: %[[C0:.*]] = arith.constant 0 : index
 // CHECK-DAG: %[[SUBVIEW:.*]] = memref.subview %[[BASE]][%[[DYN_OFFSET]], 0, 8] [1, 1, 1] [1, 1, 1] : memref<2x16x16xf32> to memref<1x1x1xf32, strided<[256, 16, 1], offset: ?>>
 // CHECK: memref.store %[[CF0]], %[[SUBVIEW]][%[[C0]], %[[C0]], %[[C0]]] {nontemporal = true} : memref<1x1x1xf32, strided<[256, 16, 1], offset: ?>>
 // CHECK: return
 func.func @test_store_nontemporal(%base : memref<2x16x16xf32>, %offset : index) -> () {
   %cf0 = arith.constant 0.0 : f32
   %c0 = arith.constant 0 : index
   %c8 = arith.constant 8 : index
   memref.store %cf0, %base[%offset, %c0, %c8] { nontemporal = true } : memref<2x16x16xf32>
   return
 }
 
 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
 }
 
 // -----
 // For this test, we made the source memref fully dynamic.
 // The gist of the check remains the same as the simple test:
 // The address computation is extracted into its own subview.
 // CHECK-LABEL: @testWithLoop(
 // CHECK-SAME: %[[BASE:[^:]*]]: memref
%sum_res1 = scf.for %iv1 = %c0 to %upper_bound1 step %c1 iter_args(%sum_iter1 = %sum_iter2) -> (f32) {
       scf.yield %sum_res0 : f32
     }
     scf.yield %sum_res1 : f32
   }
   return %sum_res2 : 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
 }
 
 // -----
 
 // Simple test: check that we extract the address computation of a ldmatrix into
 // a dedicated subview.
 // The resulting ldmatrix will loaded from with subview and have only indices set
 // to zero.
func.func @test_ldmatrix(%base : memref<4x32x32xf16, 3>,
   %loaded_val = nvgpu.ldmatrix
     %base[%offset0, %offset1, %offset2]
     {numTiles = 4 : i32, transpose = false}
       : memref<4x32x32xf16, 3> -> vector<4x2xf16>
   return %loaded_val : vector<4x2xf16>
 }
 
 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
 }
 
 // -----
 
 // Same as test_ldmatrix but with fully dynamic memref.
 
 // CHECK-DAG: #[[$A_MINUS_B_MAP:.*]] = affine_map<()[s0, s1] -> (s0 - s1)>
 // CHECK-LABEL: @test_ldmatrix(
func.func @test_ldmatrix(%base : memref<?x?x?xf16, 3>,
   %loaded_val = nvgpu.ldmatrix
     %base[%offset0, %offset1, %offset2]
     {numTiles = 4 : i32, transpose = false}
       : memref<?x?x?xf16, 3> -> vector<4x2xf16>
   return %loaded_val : vector<4x2xf16>
 }
 
 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
 }
 
 // -----
 
 // Simple test for vector.transfer_read with fully dynamic memref.
 // We also set a permutation map to make sure it is properly preserved.
 
 // CHECK-DAG: #[[$A_MINUS_B_MAP:.*]] = affine_map<()[s0, s1] -> (s0 - s1)>
func.func @test_transfer_read_op(%base : memref<?x?x?xf16>,
     %offset0 : index, %offset1: index, %offset2: index)
     -> vector<4x2xf16> {
   %cf0 = arith.constant 0.0 : f16
   %loaded_val = vector.transfer_read %base[%offset0, %offset1, %offset2], %cf0 { permutation_map = affine_map<(d0,d1,d2) -> (d2,d0)> } : memref<?x?x?xf16>, vector<4x2xf16>
   return %loaded_val : vector<4x2xf16>
 }
 
 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
 }
 
 // -----
 
 // Same as test_transfer_read_op but with tensors.
 // Right now this rewrite is not supported but we still shouldn't choke on it.
 
 // CHECK: #[[$PERMUTATION_MAP:.*]] = affine_map<(d0, d1, d2) -> (d2, d0)>
func.func @test_transfer_read_op_with_tensor(%base : tensor<?x?x?xf16>,
     %offset0 : index, %offset1: index, %offset2: index)
     -> vector<4x2xf16> {
   %cf0 = arith.constant 0.0 : f16
   %loaded_val = vector.transfer_read %base[%offset0, %offset1, %offset2], %cf0 { permutation_map = affine_map<(d0,d1,d2) -> (d2,d0)> } : tensor<?x?x?xf16>, vector<4x2xf16>
   return %loaded_val : vector<4x2xf16>
 }
 
 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
 }
 
 // -----
 
 // Simple test for vector.transfer_write with fully dynamic memref.
 // We also set a permutation map to make sure it is properly preserved.
 
 // CHECK-DAG: #[[$A_MINUS_B_MAP:.*]] = affine_map<()[s0, s1] -> (s0 - s1)>

 func.func @test_transfer_write_op(%base : memref<?x?x?xf16>,
     %offset0 : index, %offset1: index, %offset2: index) {
   %vcf0 = arith.constant dense<0.000000e+00> : vector<4x2xf16>
   vector.transfer_write %vcf0, %base[%offset0, %offset1, %offset2] { permutation_map = affine_map<(d0,d1,d2) -> (d2,d0)> } : vector<4x2xf16>, memref<?x?x?xf16>
   return
 }
 
 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
 }
 
 // -----
 
 // Check that the strides of the original memref are kept.
 // Moreover even with non-1 strides the subview should still issue [1,...]
 // strides, since this is a multiplication factor.
 

 func.func @test_transfer_write_op_with_strides(%base : memref<?x?x?xf16, strided<[329, 26, 12], offset: ?>>,
     %offset0 : index, %offset1: index, %offset2: index) {
   %vcf0 = arith.constant dense<0.000000e+00> : vector<4x2xf16>
   vector.transfer_write %vcf0, %base[%offset0, %offset1, %offset2] { permutation_map = affine_map<(d0,d1,d2) -> (d2,d0)> } : vector<4x2xf16>, memref<?x?x?xf16, strided<[329, 26, 12], offset: ?>>
   return
 }
 
 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
 }
 // -----
 
 // Same as test_transfer_write_op but with tensors.
 // Right now this rewrite is not supported but we still shouldn't choke on it.
 
 // CHECK: #[[$PERMUTATION_MAP:.*]] = affine_map<(d0, d1, d2) -> (d2, d0)>
 // CHECK-LABEL: @test_transfer_write_op_with_tensor(
 // CHECK-SAME: %[[BASE:[^:]*]]: tensor<{{[^,]*}}>,
 // CHECK-SAME: %[[DYN_OFFSET0:[^:]*]]: index,
 // CHECK-SAME: %[[DYN_OFFSET1:[^:]*]]: index,
 // CHECK-SAME: %[[DYN_OFFSET2:[^:]*]]: index)
 // CHECK-DAG: %[[VCF0:.*]] = arith.constant dense<0.0{{0*e\+00}}> : vector<4x2xf16>
 // CHECK: %[[RES:.*]] = vector.transfer_write %[[VCF0]], %[[BASE]][%[[DYN_OFFSET0]], %[[DYN_OFFSET1]], %[[DYN_OFFSET2]]] {permutation_map = #[[$PERMUTATION_MAP]]} : vector<4x2xf16>, tensor<?x?x?xf16>
 // CHECK: return %[[RES]] : tensor<?x?x?xf16>
 func.func @test_transfer_write_op_with_tensor(%base : tensor<?x?x?xf16>,
     %offset0 : index, %offset1: index, %offset2: index) -> tensor<?x?x?xf16> {
   %vcf0 = arith.constant dense<0.000000e+00> : vector<4x2xf16>
   %res = vector.transfer_write %vcf0, %base[%offset0, %offset1, %offset2] { permutation_map = affine_map<(d0,d1,d2) -> (d2,d0)> } : vector<4x2xf16>, tensor<?x?x?xf16>
   return %res : 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
 }

mlir/test/Dialect/MemRef/make-loop-independent.mlir

scf.for %i = %lb to %ub step %step {
     "test.another_use"(%alloc) : (memref<?xf32>) -> (memref<?xf32>)
 
     // CHECK: memref.store %{{.*}}, %[[subview]]
     memref.store %cst, %alloc[%c0] : memref<?xf32>
   }
   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
 }
 
 // -----
 
 // CHECK: #[[$map:.*]] = affine_map<(d0) -> (-d0 + 128)>
 // CHECK-LABEL: func @make_alloca_loop_independent_static(
 func.func @make_alloca_loop_independent_static(%step: index) {
   %cst = arith.constant 5.5 : f32
scf.for %i = %c0 to %ub step %step {
 
     // CHECK: vector.print %[[sz]]
     %dim = memref.dim %alloc, %c0 : memref<?xf32>
     vector.print %dim : 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
 }

mlir/test/Dialect/MemRef/transform-ops.mlir

scf.for %i0 = %c0 to %c16 step %c4 {
     memref.copy %1, %tmp :  memref<4xf32, affine_map<(d0)[s0] -> (d0 + s0)>> to memref<4xf32>
 
     "some_use"(%tmp) : (memref<4xf32>) ->()
   }
   return
 }
 
 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
 }
 
 // -----
 
 // CHECK-DAG: #[[$MAP0:.*]] = affine_map<(d0) -> ((d0 floordiv 4) mod 2)>
 // CHECK-DAG: #[[$MAP1:.*]] = affine_map<(d0)[s0] -> (d0 + s0)>
 
 // CHECK-LABEL: func @multi_buffer_on_affine_loop
affine.for %i0 = 0 to 16 step 4 {
     memref.copy %1, %tmp :  memref<4xf32, affine_map<(d0)[s0] -> (d0 + s0)>> to memref<4xf32>
 
     "some_use"(%tmp) : (memref<4xf32>) ->()
   }
   return
 }
 
 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
 }
 
 // -----
 
 // Trying to use multibuffer on allocs that are used in different loops
 // with none dominating the other is going to fail.
 // Check that we emit a proper error for that.
 func.func @multi_buffer_uses_with_no_loop_dominator(%in: memref<16xf32>, %cond: i1) {
scf.for %i0 = %c0 to %c16 step %c4 {
     memref.copy %1, %tmp :  memref<4xf32, affine_map<(d0)[s0] -> (d0 + s0)>> to memref<4xf32>
 
     "some_use"(%tmp) : (memref<4xf32>) ->()
   }
   return
 }
 
 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
 }
 
 // -----
 
 // Make sure the multibuffer operation is typed so that it only supports
 // memref.alloc.
 // Check that we emit an error if we try to match something else.
 func.func @multi_buffer_reject_alloca(%in: memref<16xf32>, %cond: i1) {
scf.for %i0 = %c0 to %c16 step %c4 {
     memref.copy %1, %tmp :  memref<4xf32, affine_map<(d0)[s0] -> (d0 + s0)>> to memref<4xf32>
 
     "some_use"(%tmp) : (memref<4xf32>) ->()
   }
   return
 }
 
 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
 }
 
 // -----
 
 // CHECK-DAG: #[[$MAP0:.*]] = affine_map<(d0) -> ((d0 floordiv 4) mod 2)>
 // CHECK-DAG: #[[$MAP1:.*]] = affine_map<(d0)[s0] -> (d0 + s0)>
 
 // CHECK-LABEL: func @multi_buffer_one_alloc_with_use_outside_of_loop
scf.for %i0 = %c0 to %c16 step %c4 {
     memref.copy %1, %tmp :  memref<4xf32, affine_map<(d0)[s0] -> (d0 + s0)>> to memref<4xf32>
 
     "some_use"(%tmp) : (memref<4xf32>) ->()
   }
   return
 }
 
 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
 }
 
 // -----
 
 
 // CHECK-DAG: #[[$MAP0:.*]] = affine_map<(d0) -> ((d0 floordiv 4) mod 2)>
 
 // CHECK-LABEL: func @multi_buffer
func.func @multi_buffer_no_analysis(%in: memref<16xf32>) {
   // CHECK: %[[I:.*]] = affine.apply #[[$MAP0]](%[[IV]])
   // CHECK: %[[SV:.*]] = memref.subview %[[A]][%[[I]], 0] [1, 4] [1, 1] : memref<2x4xf32> to memref<4xf32, strided<[1], offset: ?>>
     "some_write_read"(%tmp) : (memref<4xf32>) ->()
   }
   return
 }
 
 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
 }
 
 // -----
 
 // CHECK-DAG: #[[$MAP0:.*]] = affine_map<(d0) -> ((d0 floordiv 4) mod 2)>
 
 // CHECK-LABEL: func @multi_buffer_dealloc
 func.func @multi_buffer_dealloc(%in: memref<16xf32>) {
func.func @multi_buffer_dealloc(%in: memref<16xf32>) {
 
   // CHECK-NOT: memref.dealloc {{.*}} : memref<4xf32>
   // CHECK: memref.dealloc %[[A]] : memref<2x4xf32>
   memref.dealloc %tmp : memref<4xf32>
   return
 }
 
 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
 }

mlir/test/Dialect/Tensor/transform-op-make-loop-independent.mlir

scf.for %i = %lb to %ub step %step {
     } : tensor<?xf32> to tensor<?xf32>
     // CHECK: "dummy.some_use"(%[[replacement]])
     "dummy.some_use"(%p) : (tensor<?xf32>) -> ()
   }
   return
 }
 
 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
 }
 
 // -----
 
 // This is a test case where "low" padding depends on the IV.
 
 //       CHECK: #[[$map:.*]] = affine_map<()[s0, s1] -> (s0 - s1)>
 //       CHECK: #[[$map1:.*]] = affine_map<(d0)[s0, s1] -> (-d0 + s0 + s1 + 5)>
scf.for %i = %lb to %ub step %step {
     } : tensor<?xf32> to tensor<?xf32>
     // CHECK: "dummy.some_use"(%[[replacement]])
     "dummy.some_use"(%p) : (tensor<?xf32>) -> ()
   }
   return
 }
 
 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
 }
 
 // -----
 
 //       CHECK: #[[$map:.*]] = affine_map<()[s0] -> (s0 * 2 - 2)>
 // CHECK-LABEL: func @two_loops(
 func.func @two_loops(%lb: index, %ub: index, %step: index,
                      %t: tensor<?xf32>, %f: f32) {
   scf.for %i = %lb to %ub step %step {
     scf.for %j = %lb to %ub step %step {
       // CHECK: affine.apply #map()[%{{.*}}]
       %low = affine.apply affine_map<(d0, d1)[] -> (d0 + d1)> (%i, %j)[]
       %p = tensor.pad %t low[%low] high[5] {
       ^bb0(%arg1: index):
         tensor.yield %f : f32
       } : tensor<?xf32> to tensor<?xf32>
       "dummy.some_use"(%p) : (tensor<?xf32>) -> ()
     }
   }
   return
 }
 
 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
 }
 
 // -----
 
 func.func @not_enough_loops(%lb: index, %ub: index, %step: index,
                             %t: tensor<?xf32>, %f: f32) {
   scf.for %i = %lb to %ub step %step {
     scf.for %j = %lb to %ub step %step {
       %low = affine.apply affine_map<(d0, d1)[] -> (d0 + d1)> (%i, %j)[]
       // expected-note@below {{target op}}
       %p = tensor.pad %t low[%low] high[5] {
       ^bb0(%arg1: index):
         tensor.yield %f : f32
       } : tensor<?xf32> to tensor<?xf32>
       "dummy.some_use"(%p) : (tensor<?xf32>) -> ()
     }
   }
   return
 }
 
 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
 }
 
 // -----
 
 // CHECK: #[[$map:.*]] = affine_map<(d0)[s0] -> (-d0 + s0)>
 // CHECK: #[[$map1:.*]] = affine_map<()[s0, s1] -> (s0 - s1)>
 // CHECK-LABEL: func @make_empty_loop_independent(
 //  CHECK-SAME:     %[[lb:.*]]: index, %[[ub:.*]]: index, %[[step:.*]]: index)
 func.func @make_empty_loop_independent(%lb: index, %ub: index, %step: index) {
   // CHECK: scf.for %[[iv:.*]] = %[[lb]] to %[[ub]]
   scf.for %i = %lb to %ub step %step {
     // CHECK: %[[slice_sz:.*]] = affine.apply #[[$map]](%[[iv]])[%[[ub]]]
     // CHECK: %[[empty_sz:.*]] = affine.apply #[[$map1]]()[%[[ub]], %[[lb]]]
     // CHECK: %[[empty:.*]] = tensor.empty(%[[empty_sz]]) : tensor<?xf32>
     // CHECK: %[[replacement:.*]] = tensor.extract_slice %[[empty]][0] [%[[slice_sz]]] [1]
     %sz = affine.apply affine_map<(d0)[s0] -> (s0 - d0)> (%i)[%ub]
     %empty = tensor.empty(%sz) : tensor<?xf32>
     // CHECK: "dummy.some_use"(%[[replacement]])
     "dummy.some_use"(%empty) : (tensor<?xf32>) -> ()
   }
   return
 }
 
 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
 }

utils/bazel/llvm-project-overlay/mlir/BUILD.bazel

cc_library(
     deps = [
         ":AffineAnalysis",
         ":AffineDialect",
         ":AffineTransformOpsIncGen",
         ":AffineTransforms",
         ":AffineUtils",
         ":FuncDialect",
         ":IR",
-        ":PDLDialect",
         ":TransformDialect",
         ":Transforms",
         ":VectorDialect",
     ],
 )
 
 ##---------------------------------------------------------------------------##
 # AMDGPU dialect.

 
 td_library(
     name = "GPUTransformOpsTdFiles",
     srcs = [
         "include/mlir/Dialect/GPU/TransformOps/GPUTransformOps.td",
     ],
     includes = ["include"],
     deps = [
-        ":PDLDialectTdFiles",
         ":TransformDialectTdFiles",
     ],
 )
 
 gentbl_cc_library(
     name = "GPUTransformOpsIncGen",
     strip_include_prefix = "include",
     tbl_outs = [
deps = [
         ":AsmParser",
         ":ControlFlowDialect",
         ":DialectUtils",
         ":FuncDialect",
         ":GPUDialect",
         ":GPUTransformOpsIncGen",
         ":GPUTransforms",
         ":IR",
-        ":PDLDialect",
         ":Parser",
         ":SCFDialect",
         ":SideEffectInterfaces",
         ":Support",
         ":TransformDialect",
         ":TransformUtils",
         "//llvm:Support",
     ],

 
 td_library(
     name = "TensorTransformOpsTdFiles",
     srcs = [
         "include/mlir/Dialect/Tensor/TransformOps/TensorTransformOps.td",
     ],
     includes = ["include"],
     deps = [
-        ":PDLDialect",
         ":TransformDialectTdFiles",
     ],
 )
 
 gentbl_cc_library(
     name = "TensorTransformOpsIncGen",
     strip_include_prefix = "include",
     tbl_outs = [

 cc_library(
     name = "TensorTransformOps",
     srcs = glob(["lib/Dialect/Tensor/TransformOps/*.cpp"]),
     hdrs = glob(["include/mlir/Dialect/Tensor/TransformOps/*.h"]),
     includes = ["include"],
     deps = [
         ":AffineDialect",
         ":IR",
-        ":PDLDialect",
         ":SCFDialect",
         ":TensorDialect",
         ":TensorTransformOpsIncGen",
         ":TensorTransforms",
         ":TransformDialect",
         "//llvm:Support",
     ],
 )

 
 td_library(
     name = "MemRefTransformOpsTdFiles",
     srcs = [
         "include/mlir/Dialect/MemRef/TransformOps/MemRefTransformOps.td",
     ],
     includes = ["include"],
     deps = [
-        ":PDLDialect",
         ":TransformDialectTdFiles",
     ],
 )
 
 gentbl_cc_library(
     name = "MemRefTransformOpsIncGen",
     strip_include_prefix = "include",
     tbl_outs = [
deps = [
         ":AffineDialect",
         ":ArithDialect",
         ":IR",
         ":LoopLikeInterface",
         ":MemRefDialect",
         ":MemRefTransformOpsIncGen",
         ":MemRefTransforms",
         ":NVGPUDialect",
-        ":PDLDialect",
         ":SCFDialect",
         ":TransformDialect",
         ":TransformUtils",
         ":VectorDialect",
         "//llvm:Support",
     ],
 )
 

 td_library(
     name = "BufferizationTransformOpsTdFiles",
     srcs = [
         "include/mlir/Dialect/Bufferization/TransformOps/BufferizationTransformOps.td",
     ],
     includes = ["include"],
     deps = [
         ":BufferizationOpsTdFiles",
-        ":PDLDialectTdFiles",
         ":TransformDialectTdFiles",
     ],
 )
 
 gentbl_cc_library(
     name = "BufferizationTransformOpsIncGen",
     strip_include_prefix = "include",
     tbl_outs = [
cc_library(
     includes = ["include"],
     deps = [
         ":BufferizationDialect",
         ":BufferizationEnumsIncGen",
         ":BufferizationTransformOpsIncGen",
         ":BufferizationTransforms",
         ":IR",
         ":MemRefDialect",
-        ":PDLDialect",
         ":Parser",
         ":SideEffectInterfaces",
         ":TensorDialect",
         ":TransformDialect",
         "//llvm:Support",
     ],
 )