[mlir] switch transform dialect ops to use TransformTypeInterface

[ftynse]

Summary

Use the recently introduced TransformTypeInterface instead of hardcoding
the PDLOperationType. This will allow the operations to use more
specific transform types to express pre/post-conditions in the future.
It requires the syntax and Python op construction API to be updated.
Dialect extensions will be switched separately.

[ftynse]

Rebase.

[ftynse]

Rebase.

mlir/include/mlir/Dialect/Transform/IR/TransformOps.td

 //===- TransformOps.td - Transform dialect operations ------*- 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 MLIR_DIALECT_TRANSFORM_IR_TRANSFORMOPS
 #define MLIR_DIALECT_TRANSFORM_IR_TRANSFORMOPS
 
 include "mlir/Interfaces/CastInterfaces.td"
 include "mlir/Interfaces/ControlFlowInterfaces.td"
+include "mlir/Interfaces/InferTypeOpInterface.td"
 include "mlir/IR/OpAsmInterface.td"
 include "mlir/IR/SymbolInterfaces.td"
-include "mlir/Dialect/PDL/IR/PDLTypes.td"
 include "mlir/Dialect/Transform/IR/TransformAttrs.td"
 include "mlir/Dialect/Transform/IR/TransformDialect.td"
 include "mlir/Dialect/Transform/IR/TransformEffects.td"
 include "mlir/Dialect/Transform/IR/TransformInterfaces.td"
 
 def AlternativesOp : TransformDialectOp<"alternatives",
     [DeclareOpInterfaceMethods<RegionBranchOpInterface,
         ["getSuccessorEntryOperands", "getSuccessorRegions",
let description = [{
     ^bb0(%arg0: !pdl.operation):
       // Otherwise, the second alternative is tried and it always succeeds by
       // returning the original handle.
       transform.yield %arg0 : !pdl.operation
     }
     ```
   }];
 
-  let arguments = (ins Optional<PDL_Operation>:$scope);
-  let results = (outs Variadic<AnyType>:$results);
+  let arguments = (ins Optional<TransformTypeInterface>:$scope);
+  let results = (outs Variadic<TransformTypeInterface>:$results);
   let regions = (region VariadicRegion<SizedRegion<1>>:$alternatives);
 
   let assemblyFormat =
-    "($scope^)? (`->` type($results)^)? attr-dict-with-keyword regions";
+    "($scope^ `:` type($scope))? (`->` type($results)^)? "
+    "attr-dict-with-keyword regions";
   let hasVerifier = 1;
 }
 
 def CastOp : TransformDialectOp<"cast",
     [TransformOpInterface, TransformEachOpTrait,
      DeclareOpInterfaceMethods<CastOpInterface>,
      DeclareOpInterfaceMethods<MemoryEffectsOpInterface>]> {
-  // TODO: temporarily fallback support for casting from PDL_Operation type.
-  let arguments = (ins AnyType:$input);
-  let results = (outs AnyType:$output);
+  let arguments = (ins TransformTypeInterface:$input);
+  let results = (outs TransformTypeInterface:$output);
   let assemblyFormat = "$input attr-dict `:` type($input) `to` type($output)";
 
   let extraClassDeclaration = [{
     ::mlir::DiagnosedSilenceableFailure applyToOne(
       ::mlir::Operation *target,
       ::llvm::SmallVectorImpl<::mlir::Operation *> &results,
       ::mlir::transform::TransformState &state);
   }];
let description = [{
     invalid state, i.e., this operation offers no transformation rollback
     capabilities.
 
     This op generates as many handles as the terminating YieldOp has operands.
     For each result, the payload ops of the corresponding YieldOp operand are
     merged and mapped to the same resulting handle.
   }];
 
-  let arguments = (ins PDL_Operation:$target);
-  let results = (outs Variadic<PDL_Operation>:$results);
+  let arguments = (ins TransformTypeInterface:$target);
+  let results = (outs Variadic<TransformTypeInterface>:$results);
   let regions = (region SizedRegion<1>:$body);
-  let assemblyFormat = "$target (`->` type($results)^)? $body attr-dict";
+  let assemblyFormat = 
+    "$target `:` type($target) (`->` type($results)^)? $body attr-dict";
   let hasVerifier = 1;
 
   let extraClassDeclaration = [{
     /// Allow the dialect prefix to be omitted.
     static StringRef getDefaultDialect() { return "transform"; }
 
     BlockArgument getIterationVariable() {
       return getBody().front().getArgument(0);
let description = [{
     with the operand have the same parent) for which the ancestor will only be
     listed once for the first time it occurs. For example, given the list
     "(childof(A), childof(B), childof(B), childof(A), childof(B))", the
     resulting list will be just "(A, B)". Note that no other semantic ordering
     is applied, e.g., "B" may itself be a parent of "A". This may have an impact
     on the further transformation applied to the handle produced here.
   }];
 
-  let arguments = (ins PDL_Operation:$target);
-  let results = (outs PDL_Operation:$parent);
-  let assemblyFormat = "$target attr-dict";
+  let arguments = (ins TransformTypeInterface:$target);
+  let results = (outs TransformTypeInterface:$parent);
+  let assemblyFormat =
+    "$target attr-dict `:` functional-type(operands, results)";
 }
 
 def GetProducerOfOperand : TransformDialectOp<"get_producer_of_operand",
     [DeclareOpInterfaceMethods<TransformOpInterface>,
      NavigationTransformOpTrait, MemoryEffectsOpInterface]> {
   let summary = "Get handle to the producer of this operation's operand number";
   let description = [{
     The handle defined by this Transform op corresponds to operation that
     produces the SSA value defined by the `target` and `operand_number`
     arguments. If the origin of the SSA value is not an operations (i.e. it is
     a block argument), the transform silently fails.
     The return handle points to only the subset of successfully produced
     computational operations, which can be empty.
   }];
 
-  let arguments = (ins PDL_Operation:$target,
+  let arguments = (ins TransformTypeInterface:$target,
                        I64Attr:$operand_number);
-  let results = (outs PDL_Operation:$parent);
-  let assemblyFormat = "$target `[` $operand_number `]` attr-dict";
+  let results = (outs TransformTypeInterface:$parent);
+  let assemblyFormat = "$target `[` $operand_number `]` attr-dict `:` "
+                       "functional-type(operands, results)";
 }
 
 def MergeHandlesOp : TransformDialectOp<"merge_handles",
     [DeclareOpInterfaceMethods<TransformOpInterface>,
-     DeclareOpInterfaceMethods<MemoryEffectsOpInterface>]> {
+     DeclareOpInterfaceMethods<MemoryEffectsOpInterface>,
+     SameOperandsAndResultType]> {
   let summary = "Merges handles into one pointing to the union of payload ops";
   let description = [{
     Creates a new Transform IR handle value that points to the same Payload IR
     operations as the operand handles. The Payload IR operations are listed
     in the same order as they are in the operand handles, grouped by operand
     handle, e.g., all Payload IR operations associated with the first handle
     come first, then all Payload IR operations associated with the second handle
     and so on. If `deduplicate` is set, do not add the given Payload IR
     operation more than once to the final list regardless of it coming from the
     same or different handles. Consumes the operands and produces a new handle.
   }];
 
-  let arguments = (ins Variadic<PDL_Operation>:$handles,
+  let arguments = (ins Variadic<TransformTypeInterface>:$handles,
                        UnitAttr:$deduplicate);
-  let results = (outs PDL_Operation:$result);
-  let assemblyFormat = "($deduplicate^)? $handles attr-dict";
+  let results = (outs TransformTypeInterface:$result);
+  let assemblyFormat = "($deduplicate^)? $handles attr-dict `:` type($result)";
   let hasFolder = 1;
 }
 
 def SplitHandlesOp : TransformDialectOp<"split_handles",
     [FunctionalStyleTransformOpTrait,
      DeclareOpInterfaceMethods<TransformOpInterface>,
      DeclareOpInterfaceMethods<MemoryEffectsOpInterface>]> {
   let summary = "Splits handles from a union of payload ops to a list";
   let description = [{
     Creates `num_result_handles` transform IR handles extracted from the 
     `handle` operand. The resulting Payload IR operation handles are listed
     in the same order as the operations appear in the source `handle`.
     This is useful for ensuring a statically known number of operations are
     tracked by the source `handle` and to extract them into individual handles
     that can be further manipulated in isolation.
 
     This operation succeeds and returns `num_result_handles` if the statically
     specified `num_result_handles` corresponds to the dynamic number of 
     operations contained in the source `handle`. Otherwise it silently fails.
   }];
 
-  let arguments = (ins PDL_Operation:$handle,
+  let arguments = (ins TransformTypeInterface:$handle,
                        I64Attr:$num_result_handles);
-  let results = (outs Variadic<PDL_Operation>:$results);
+  let results = (outs Variadic<TransformTypeInterface>:$results);
   let assemblyFormat = [{
     $handle `in` `[` $num_result_handles `]` 
-    custom<StaticNumPDLResults>(type($results), ref($num_result_handles))
-    attr-dict
+    attr-dict `:` functional-type(operands, results)
   }];
 }
 
 def PDLMatchOp : TransformDialectOp<"pdl_match",
     [DeclareOpInterfaceMethods<TransformOpInterface>,
      DeclareOpInterfaceMethods<MemoryEffectsOpInterface>]> {
   let summary = "Finds ops that match the named PDL pattern";
   let description = [{
let description = [{
     of the matched ops.
 
     The transformation is considered successful regardless of whether some
     Payload IR ops actually matched the pattern and only fails if the pattern
     could not be looked up or compiled.
   }];
 
   let arguments = (ins
-    Arg<PDL_Operation, "Payload IR scope to match within">:$root, 
+    Arg<TransformTypeInterface, "Payload IR scope to match within">:$root, 
     SymbolRefAttr:$pattern_name);
   let results = (outs
-    Res<PDL_Operation, "Handle to the matched Payload IR ops">:$matched);
+    Res<TransformTypeInterface, "Handle to the matched Payload IR ops">:$matched);
 
-  let assemblyFormat = "$pattern_name `in` $root attr-dict";
+  let assemblyFormat = "$pattern_name `in` $root attr-dict `:` "
+                       "functional-type(operands, results)";
 }
 
 def ReplicateOp : TransformDialectOp<"replicate",
     [DeclareOpInterfaceMethods<TransformOpInterface>,
-     DeclareOpInterfaceMethods<MemoryEffectsOpInterface>]> {
+     DeclareOpInterfaceMethods<MemoryEffectsOpInterface>,
+     AllTypesMatch<["handles", "replicated"]>]> {
   let summary = "Lists payload ops multiple times in the new handle";
   let description = [{
     Produces a new handle associated with a list of payload IR ops that is
     computed by repeating the list of payload IR ops associated with the
     operand handle as many times as the "pattern" handle has associated
     operations. For example, if pattern is associated with [op1, op2] and the
     operand handle is associated with [op3, op4, op5], the resulting handle
     will be associated with [op3, op4, op5, op3, op4, op5].
let description = [{
     an operation that consumes the handle. Handle consumption often indicates
     that the associated payload IR ops are destroyed, so having the same op
     listed more than once will lead to double-free. Single-operand
     MergeHandlesOp may be used to deduplicate the associated list of payload IR
     ops when necessary. Furthermore, a combination of ReplicateOp and
     MergeHandlesOp can be used to construct arbitrary lists with repetitions.
   }];
 
-  let arguments = (ins PDL_Operation:$pattern,
-                       Variadic<PDL_Operation>:$handles);
-  let results = (outs Variadic<PDL_Operation>:$replicated);
-  let assemblyFormat =
-    "`num` `(` $pattern `)` $handles "
-    "custom<PDLOpTypedResults>(type($replicated), ref($handles)) attr-dict";
+  let arguments = (ins TransformTypeInterface:$pattern,
+                       Variadic<TransformTypeInterface>:$handles);
+  let results = (outs Variadic<TransformTypeInterface>:$replicated);
+  let assemblyFormat = "`num` `(` $pattern `)` $handles attr-dict `:` "
+                       "type($pattern) `,` type($handles)";
 }
 
 def SequenceOp : TransformDialectOp<"sequence",
     [DeclareOpInterfaceMethods<RegionBranchOpInterface,
         ["getSuccessorEntryOperands", "getSuccessorRegions",
          "getRegionInvocationBounds"]>,
      DeclareOpInterfaceMethods<TransformOpInterface>,
      DeclareOpInterfaceMethods<MemoryEffectsOpInterface>,
let description = [{
     The entry block of this operation has a single argument that maps to either
     the operand if provided or the top-level container operation of the payload
     IR, typically the root operation of the pass interpreting the transform
     dialect. Operand omission is only allowed for sequences not contained in
     another sequence.
   }];
 
   let arguments = (ins FailurePropagationMode:$failure_propagation_mode,
-                       Optional<PDL_Operation>:$root);
-  let results = (outs Variadic<AnyType>:$results);
+                       Optional<TransformTypeInterface>:$root);
+  let results = (outs Variadic<TransformTypeInterface>:$results);
   let regions = (region SizedRegion<1>:$body);
 
   let assemblyFormat =
-    "($root^)? `failures` `(` $failure_propagation_mode `)` attr-dict-with-keyword regions (`:` type($results)^)?";
+    "($root^ `:` type($root))? (`->` type($results)^)? `failures` `(` "
+    "$failure_propagation_mode `)` attr-dict-with-keyword regions";
 
   let extraClassDeclaration = [{
     /// Allow the dialect prefix to be omitted.
     static StringRef getDefaultDialect() { return "transform"; }
   }];
 
   let hasVerifier = 1;
 }
let description = [{
     This op is expected to contain `pdl.pattern` operations and exactly one
     another Transform dialect operation that gets executed with all patterns
     available. This op is a possible top-level Transform IR op, the argument of
     its entry block corresponds to either the root op of the payload IR or the
     ops associated with its operand when provided.
   }];
 
   let arguments = (ins
-    Arg<Optional<PDL_Operation>, "Root operation of the Payload IR",
+    Arg<Optional<TransformTypeInterface>, "Root operation of the Payload IR",
         [TransformMappingRead]>:$root);
   let regions = (region SizedRegion<1>:$body);
-  let assemblyFormat = "($root^)? attr-dict-with-keyword regions";
+  let assemblyFormat = "($root^ `:` type($root))? attr-dict-with-keyword regions";
 
   let hasVerifier = 1;
 
   let extraClassDeclaration = [{
     /// Allow the dialect prefix to be omitted.
     static StringRef getDefaultDialect() { return "transform"; }
   }];
 }
 
 def YieldOp : TransformDialectOp<"yield", [Terminator]> {
   let summary = "Yields operation handles from a transform IR region";
   let description = [{
     This terminator operation yields operation handles from regions of the
     transform IR ops back to the containing op. It is not itself associated with
     any transformation on the payload IR and is used for flow purposes only.
   }];
 
   let arguments = (ins 
-    Arg<Variadic<AnyType>, "Operation handles yielded back to the parent",
+    Arg<Variadic<TransformTypeInterface>, "Operation handles yielded back to the parent",
         [TransformMappingRead]>:$operands);
   let assemblyFormat = "operands attr-dict (`:` type($operands)^)?";
 
   let builders = [
     OpBuilder<(ins), [{
       return build($_builder, $_state, ::mlir::ValueRange());
     }]>
   ];
 }
 
 #endif // MLIR_DIALECT_TRANSFORM_IR_TRANSFORMOPS

mlir/lib/Dialect/Transform/IR/TransformDialect.cpp

 
 void transform::detail::checkImplementsTransformTypeInterface(
     TypeID typeID, MLIRContext *context) {
   const auto &abstractType = AbstractType::lookup(typeID, context);
   assert(abstractType.hasInterface(TransformTypeInterface::getInterfaceID()));
 }
 #endif // NDEBUG
 
+namespace {
+struct PDLOperationTypeTransformTypeInterfaceImpl
+    : public transform::TransformTypeInterface::ExternalModel<
+          PDLOperationTypeTransformTypeInterfaceImpl, pdl::OperationType> {
+  DiagnosedSilenceableFailure
+  checkPayload(Type type, Location loc, ArrayRef<Operation *> payload) const {
+    return DiagnosedSilenceableFailure::success();
+  }
+};
+} // namespace
+
 void transform::TransformDialect::initialize() {
   // Using the checked versions to enable the same assertions as for the ops
   // from extensions.
   addOperationsChecked<
 #define GET_OP_LIST
 #include "mlir/Dialect/Transform/IR/TransformOps.cpp.inc"
       >();
   addTypesChecked<
 #define GET_TYPEDEF_LIST
 #include "mlir/Dialect/Transform/IR/TransformTypes.cpp.inc"
       >();
+
+  pdl::OperationType::attachInterface<
+      PDLOperationTypeTransformTypeInterfaceImpl>(*getContext());
 }
 
 void transform::TransformDialect::mergeInPDLMatchHooks(
     llvm::StringMap<PDLConstraintFunction> &&constraintFns) {
   // Steal the constraint functions form the given map.
   for (auto &it : constraintFns)
     pdlMatchHooks.registerConstraintFunction(it.getKey(), std::move(it.second));
 }

mlir/lib/Dialect/Transform/IR/TransformInterfaces.cpp

 //===- TransformInterfaces.cpp - Transform Dialect Interfaces -------------===//
 //
 // 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/Transform/IR/TransformInterfaces.h"
-#include "mlir/Dialect/PDL/IR/PDLTypes.h"
 #include "mlir/Dialect/Transform/IR/TransformTypes.h"
 #include "mlir/IR/Diagnostics.h"
 #include "mlir/IR/Operation.h"
 #include "llvm/ADT/STLExtras.h"
 #include "llvm/Support/Debug.h"
 
 #define DEBUG_TYPE "transform-dialect"
 #define DEBUG_PRINT_AFTER_ALL "transform-dialect-print-top-level-after-all"
transform::TransformState::setPayloadOps(Value value,
                                          ArrayRef<Operation *> targets) {
   assert(value != kTopLevelValue &&
          "attempting to reset the transformation root");
 
   // TODO: this may go now
   if (value.use_empty())
     return success();
 
-  if (auto iface = value.getType().dyn_cast<TransformTypeInterface>()) {
-    DiagnosedSilenceableFailure result =
-        iface.checkPayload(value.getLoc(), targets);
-    if (failed(result.checkAndReport()))
-      return failure();
-  }
+  auto iface = value.getType().cast<TransformTypeInterface>();
 
   // Setting new payload for the value without cleaning it first is a misuse of
   // the API, assert here.
   SmallVector<Operation *> storedTargets(targets.begin(), targets.end());
   Mappings &mappings = getMapping(value);
   bool inserted =
       mappings.direct.insert({value, std::move(storedTargets)}).second;
   assert(inserted && "value is already associated with another list");
LogicalResult transform::TransformState::updatePayloadOps(
   for (Operation *op : association) {
     dropReverseMapping(mappings, op, value);
     if (Operation *updatedOp = callback(op)) {
       updated.push_back(updatedOp);
       mappings.reverse[updatedOp].push_back(value);
     }
   }
 
-  if (auto iface = value.getType().dyn_cast<TransformTypeInterface>()) {
-    DiagnosedSilenceableFailure result =
-        iface.checkPayload(value.getLoc(), updated);
-    if (failed(result.checkAndReport()))
-      return failure();
-  }
+  auto iface = value.getType().cast<TransformTypeInterface>();
 
   std::swap(association, updated);
   return success();
 }
 
 void transform::TransformState::recordHandleInvalidationOne(
     OpOperand &handle, Operation *payloadOp, Value otherHandle) {
   ArrayRef<Operation *> potentialAncestors = getPayloadOps(handle.get());
if (op->getNumRegions() < 1)
     return op->emitOpError() << "expects at least one region";
 
   Region *bodyRegion = &op->getRegion(0);
   if (!llvm::hasNItems(*bodyRegion, 1))
     return op->emitOpError() << "expects a single-block region";
 
   Block *body = &bodyRegion->front();
   if (body->getNumArguments() != 1 ||
-      !body->getArgumentTypes()[0].isa<pdl::OperationType>()) {
-    return op->emitOpError()
-           << "expects the entry block to have one argument of type "
-           << pdl::OperationType::get(op->getContext());
+      !body->getArgumentTypes()[0].isa<TransformTypeInterface>()) {
+    return op->emitOpError() << "expects the entry block to have one argument "
+                                "of type implementing TransformTypeInterface";
   }
 
   if (auto *parent =
           op->getParentWithTrait<PossibleTopLevelTransformOpTrait>()) {
     if (op->getNumOperands() == 0) {
       InFlightDiagnostic diag =
           op->emitOpError()
           << "expects the root operation to be provided for a nested op";

mlir/lib/Dialect/Transform/IR/TransformOps.cpp

 //===- TransformDialect.cpp - Transform dialect operations ----------------===//
 //
 // 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/Transform/IR/TransformOps.h"
 #include "mlir/Dialect/PDL/IR/PDLOps.h"
 #include "mlir/Dialect/Transform/IR/TransformDialect.h"
 #include "mlir/Dialect/Transform/IR/TransformInterfaces.h"
+#include "mlir/Dialect/Transform/IR/TransformTypes.h"
 #include "mlir/IR/OpImplementation.h"
 #include "mlir/IR/PatternMatch.h"
 #include "mlir/Interfaces/ControlFlowInterfaces.h"
 #include "mlir/Rewrite/FrozenRewritePatternSet.h"
 #include "mlir/Rewrite/PatternApplicator.h"
 #include "llvm/ADT/STLExtras.h"
 #include "llvm/ADT/ScopeExit.h"
 #include "llvm/Support/Debug.h"
 
 #define DEBUG_TYPE "transform-dialect"
 #define DBGS() (llvm::dbgs() << "[" DEBUG_TYPE "] ")
 
 using namespace mlir;
 
-/// Custom parser for ReplicateOp.
-static ParseResult parsePDLOpTypedResults(
-    OpAsmParser &parser, SmallVectorImpl<Type> &types,
-    const SmallVectorImpl<OpAsmParser::UnresolvedOperand> &handles) {
-  types.resize(handles.size(), pdl::OperationType::get(parser.getContext()));
-  return success();
-}
-
-/// Custom printer for ReplicateOp.
-static void printPDLOpTypedResults(OpAsmPrinter &, Operation *, TypeRange,
-                                   ValueRange) {}
-
-/// Custom parser for SplitHandlesOp.
-static ParseResult parseStaticNumPDLResults(OpAsmParser &parser,
-                                            SmallVectorImpl<Type> &types,
-                                            IntegerAttr numHandlesAttr) {
-  types.resize(numHandlesAttr.getInt(),
-               pdl::OperationType::get(parser.getContext()));
-  return success();
-}
-
-/// Custom printer for SplitHandlesOp.
-static void printStaticNumPDLResults(OpAsmPrinter &, Operation *, TypeRange,
-                                     IntegerAttr) {}
-
 #define GET_OP_CLASSES
 #include "mlir/Dialect/Transform/IR/TransformOps.cpp.inc"
 
 //===----------------------------------------------------------------------===//
 // PatternApplicatorExtension
 //===----------------------------------------------------------------------===//
 
 namespace {
for (Region &reg : getAlternatives()) {
     }
   }
   return emitSilenceableError() << "all alternatives failed";
 }
 
 LogicalResult transform::AlternativesOp::verify() {
   for (Region &alternative : getAlternatives()) {
     Block &block = alternative.front();
-    if (block.getNumArguments() != 1 ||
-        !block.getArgument(0).getType().isa<pdl::OperationType>()) {
-      return emitOpError()
-             << "expects region blocks to have one operand of type "
-             << pdl::OperationType::get(getContext());
-    }
-
     Operation *terminator = block.getTerminator();
     if (terminator->getOperands().getTypes() != getResults().getTypes()) {
       InFlightDiagnostic diag = emitOpError()
                                 << "expects terminator operands to have the "
                                    "same type as results of the operation";
       diag.attachNote(terminator->getLoc()) << "terminator";
       return diag;
     }

 }
 
 LogicalResult transform::ForeachOp::verify() {
   auto yieldOp = getYieldOp();
   if (getNumResults() != yieldOp.getNumOperands())
     return emitOpError() << "expects the same number of results as the "
                             "terminator has operands";
   for (Value v : yieldOp.getOperands())
-    if (!v.getType().isa<pdl::OperationType>())
-      return yieldOp->emitOpError("expects only PDL_Operation operands");
+    if (!v.getType().isa<TransformTypeInterface>())
+      return yieldOp->emitOpError(
+          "expects operands to have types implementing TransformTypeInterface");
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // GetClosestIsolatedParentOp
 //===----------------------------------------------------------------------===//
 
 DiagnosedSilenceableFailure transform::GetClosestIsolatedParentOp::apply(

mlir/python/mlir/dialects/_transform_ops_ext.py

 #  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
 
 try:
   from ..ir import *
   from ._ods_common import get_op_result_or_value as _get_op_result_or_value, get_op_results_or_values as _get_op_results_or_values
-  from ..dialects import pdl
 except ImportError as e:
   raise RuntimeError("Error loading imports from extension module") from e
 
 from argparse import SUPPRESS
 from typing import Optional, overload, Sequence, Union
 
 
 def _get_symbol_ref_attr(value: Union[Attribute, str]):
   if isinstance(value, Attribute):
     return value
   return FlatSymbolRefAttr.get(value)
 
 
 class GetClosestIsolatedParentOp:
 
-  def __init__(self, target: Union[Operation, Value], *, loc=None, ip=None):
+  def __init__(self, result_type: Type, target: Union[Operation, Value], *, loc=None, ip=None):
     super().__init__(
-        pdl.OperationType.get(),
+        result_type,
         _get_op_result_or_value(target),
         loc=loc,
         ip=ip)
 
 
 class MergeHandlesOp:
 
   def __init__(self,
                handles: Sequence[Union[Operation, Value]],
                *,
                deduplicate: bool = False,
                loc=None,
                ip=None):
     super().__init__(
-        pdl.OperationType.get(), [_get_op_result_or_value(h) for h in handles],
+        [_get_op_result_or_value(h) for h in handles],
         deduplicate=deduplicate,
         loc=loc,
         ip=ip)
 
 
 class PDLMatchOp:
 
   def __init__(self,
+               result_type: Type,
                target: Union[Operation, Value],
                pattern_name: Union[Attribute, str],
                *,
                loc=None,
                ip=None):
     super().__init__(
-        pdl.OperationType.get(),
+        result_type,
         _get_op_result_or_value(target),
         _get_symbol_ref_attr(pattern_name),
         loc=loc,
         ip=ip)
 
 
 class ReplicateOp:
 
   def __init__(self,
                pattern: Union[Operation, Value],
                handles: Sequence[Union[Operation, Value]],
                *,
                loc=None,
                ip=None):
     super().__init__(
-        [pdl.OperationType.get()] * len(handles),
+        [_get_op_result_or_value(h).type for h in handles],
         _get_op_result_or_value(pattern),
         [_get_op_result_or_value(h) for h in handles],
         loc=loc,
         ip=ip)
 
 
 class SequenceOp:
 
-  @overload
-  def __init__(self, failure_propagation_mode,
-               resultsOrRoot: Sequence[Type],
-               optionalRoot: Optional[Union[Operation, Value]]):
-    ...
-
-  @overload
-  def __init__(self, failure_propagation_mode,
-               resultsOrRoot: Optional[Union[Operation,
-                                             Value]], optionalRoot: NoneType):
-    ...
-
-  def __init__(self, failure_propagation_mode, resultsOrRoot=None, optionalRoot=None):
-    results = resultsOrRoot if isinstance(resultsOrRoot, Sequence) else []
-    root = (
-        resultsOrRoot
-        if not isinstance(resultsOrRoot, Sequence) else optionalRoot)
-    root = _get_op_result_or_value(root) if root else None
+  def __init__(self, failure_propagation_mode, results: Sequence[Type],
+               target: Union[Operation, Value, Type]):
+    root = _get_op_result_or_value(target) if isinstance(
+        target, (Operation, Value)) else None
+    root_type = root.type if not isinstance(target, Type) else target
     if not isinstance(failure_propagation_mode, Attribute):
       failure_propagation_mode_attr = IntegerAttr.get(
           IntegerType.get_signless(32), failure_propagation_mode._as_int())
     else:
       failure_propagation_mode = failure_propagation_mode
     super().__init__(results_=results,
                      failure_propagation_mode=failure_propagation_mode_attr,
                      root=root)
-    self.regions[0].blocks.append(pdl.OperationType.get())
+    self.regions[0].blocks.append(root_type)
 
   @property
   def body(self) -> Block:
     return self.regions[0].blocks[0]
 
   @property
   def bodyTarget(self) -> Value:
     return self.body.arguments[0]
 
 
 class WithPDLPatternsOp:
 
   def __init__(self,
-               target: Optional[Union[Operation, Value]] = None,
+               target: Union[Operation, Value, Type],
                *,
                loc=None,
                ip=None):
+    root = _get_op_result_or_value(target) if not isinstance(target,
+                                                             Type) else None
+    root_type = target if isinstance(target, Type) else root.type
     super().__init__(
-        root=_get_op_result_or_value(target) if target else None,
+        root=root,
         loc=loc,
         ip=ip)
-    self.regions[0].blocks.append(pdl.OperationType.get())
+    self.regions[0].blocks.append(root_type)
 
   @property
   def body(self) -> Block:
     return self.regions[0].blocks[0]
 
   @property
   def bodyTarget(self) -> Value:
     return self.body.arguments[0]

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.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
     ^bb0(%arg1: !pdl.operation):
       %0 = transform.structured.match ops{["func.func"]} in %arg1
       transform.bufferization.one_shot_bufferize %0
           {target_is_module = false}
   }
 }
 
 // CHECK-LABEL: func @test_function(

 }
 
 // -----
 
 // Test analysis of One-Shot Bufferize only.
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
     ^bb0(%arg1: !pdl.operation):
       %0 = transform.structured.match ops{["func.func"]} in %arg1
       transform.bufferization.one_shot_bufferize %0
           {target_is_module = false, test_analysis_only = true}
   }
 }
 
 // CHECK-LABEL: func @test_function_analysis(

 
 // -----
 
 // Test One-Shot Bufferize transform failure with an unknown op. This would be
 // allowed with `allow_unknown_ops`.
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
     ^bb0(%arg1: !pdl.operation):
       %0 = transform.structured.match ops{["func.func"]} in %arg1
       // expected-error @+1 {{bufferization failed}}
       transform.bufferization.one_shot_bufferize %0 {target_is_module = false}
   }
 }
 
 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>
 }
 
 // -----
 
 // Test One-Shot Bufferize transform failure with a module op.
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
     ^bb0(%arg1: !pdl.operation):
       // %arg1 is the module
       transform.bufferization.one_shot_bufferize %arg1
   }
 }
 
 module {
   // CHECK-LABEL: func @test_function(

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

%name2 = gpu.launch async[%stream] blocks(%arg3, %arg4, %arg5) in (%arg9 = %one, %arg10 = %one, %arg11 = %one)
      }  {thread_dim_mapping = [1, 0, 2]}
     gpu.terminator
   }
 
   return %y : memref<2 x 32 x f32>
 }
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0
     // expected-error @below {{Trying to launch a GPU kernel with gridDim = (1, 1, 1) blockDim = (1200, 9, 1). It is larger than the limits.}}
     // expected-note @below {{"blockDim" is very large}}
     transform.gpu.map_nested_foreach_to_threads %funcop { blockDim = [1200, 9, 1] }
   }
 }
 
%name2 = gpu.launch async[%stream] blocks(%arg3, %arg4, %arg5) in (%arg9 = %one, %arg10 = %one, %arg11 = %one)
      }  {thread_dim_mapping = [1, 0, 2]}
     gpu.terminator
   }
 
   return %y : memref<2 x 32 x f32>
 }
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0
     // expected-error @below {{The requested GPU threads are fewer than the number of loop trip counts. Try to tile scf.foreach_thread before mapping or set small blockDim.}}
     transform.gpu.map_nested_foreach_to_threads %funcop { blockDim = [128, 4, 1] }
   }
 }
 
 // -----
%name = gpu.launch async[%stream] blocks(%arg3, %arg4, %arg5) in (%arg9 = %one, %arg10 = %one, %arg11 = %one)
      }  {thread_dim_mapping = [1, 0, 2]}
     gpu.terminator
   }
   return %y : memref<2 x 32 x f32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0
     // expected-error @below {{unsupported dynamic blockdim size}}
     transform.gpu.map_nested_foreach_to_threads %funcop { blockDim = [128, 4, 1] }
   }
 }
 
 // -----
%name = gpu.launch async[%stream] blocks(%arg3, %arg4, %arg5) in (%arg9 = %one, %arg10 = %one, %arg11 = %one)
      }  {thread_dim_mapping = [1, 0, 2]}
     gpu.terminator
   }
   return %y : memref<2x32x32x32xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0
     // expected-error @below {{scf.foreach_thread with rank > 3 does not lower to gpu.thread_id}}
     transform.gpu.map_nested_foreach_to_threads %funcop { blockDim = [128, 4, 1] }
   }
 }
 
 // -----
 
 func.func @map_nested_foreach_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.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %matmul = transform.structured.match ops{["linalg.matmul"]} in %arg0
     %foreach, %tiled = transform.structured.tile_to_foreach_thread_op %matmul num_threads [10, 20, 30]
     %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0
     // expected-error @below {{only bufferized scf.foreach_thread lowers to gpu.thread_id}}    
     transform.gpu.map_nested_foreach_to_threads %funcop { blockDim = [128, 4, 1] }
   }
 }

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

func.func @saxpy2dblock(%x: !type, %y: !type, %t: !type1d, %alpha : f32, %stream : !gpu.async.token) -> !type {
   %c7 = arith.constant 7 : index
   %one = arith.constant 1 : index
 //      CHECK:   gpu.launch
 //      CHECK:   %[[BLKX:.*]] = gpu.block_id  x
 //      CHECK:   %[[BLKY:.*]] = gpu.block_id  y
 //      CHECK:   memref.load %[[ARGX]][%[[BLKX]], %[[BLKY]]]
 //      CHECK:   memref.load %[[ARGY]][%[[BLKX]], %[[BLKY]]]
   %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) 
+            threads(%arg6, %arg7, %arg8) in (%arg12 = %one, %arg13 = %one, %arg14 = %one)
   {
     scf.foreach_thread (%i, %j) in (%c7, %c9) {
-        %4 = memref.load %x[%i, %j] : !type        
+        %4 = memref.load %x[%i, %j] : !type
         %5 = memref.load %y[%i, %j] : !type
         %6 = math.fma %alpha, %4, %5 : f32
         memref.store %6, %y[%i, %j] : !type
-     }  {thread_dim_mapping = [0, 1, 2]}               
+     }  {thread_dim_mapping = [0, 1, 2]}
     gpu.terminator
-  }  
+  }
   return %y : !type
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0
     transform.gpu.map_foreach_to_blocks %funcop { blockDim = [12, 9, 1]}
   }
 }
 
 // -----
 

 //      CHECK:   arith.cmpi ult, %[[TIDY]], %[[C7]] : index
 //      CHECK:   memref.load %[[ARGX]][%[[TIDY]], %[[TIDX]]]
 //      CHECK:   memref.load %[[ARGY]][%[[TIDY]], %[[TIDX]]]
 //      CHECK:   gpu.barrier
 //      CHECK:   arith.cmpi ult, %[[TIDY]], %[[C1]] : index
 //      CHECK:   memref.load %[[ARGT]][%[[TIDX]]]
 //      CHECK:   gpu.barrier
   %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) 
+            threads(%arg6, %arg7, %arg8) in (%arg12 = %one, %arg13 = %one, %arg14 = %one)
   {
     scf.foreach_thread (%i, %j) in (%c7, %c9) {
-        %4 = memref.load %x[%i, %j] : !type        
+        %4 = memref.load %x[%i, %j] : !type
         %5 = memref.load %y[%i, %j] : !type
         %6 = math.fma %alpha, %4, %5 : f32
         memref.store %6, %y[%i, %j] : !type
-     }  {thread_dim_mapping = [1, 0, 2]}    
+     }  {thread_dim_mapping = [1, 0, 2]}
      scf.foreach_thread (%i) in (%c12) {
-        %7 = memref.load %t[%i] : !type1d        
+        %7 = memref.load %t[%i] : !type1d
         %8 = arith.addf %alpha, %7 : f32
         memref.store %8, %t[%i] : !type1d
      }  {thread_dim_mapping = [0, 1, 2]}
     gpu.terminator
   }
   return %y : !type
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0
     transform.gpu.map_nested_foreach_to_threads %funcop { blockDim = [12, 9, 1] }
   }
 }
 
 // -----
 
 !type4d = memref<32x64x4x32xf32>
 
 // CHECK-LABEL: func.func @saxpy4d(
 // CHECK-SAME:    %[[ARGX:[0-9a-z]+]]: memref<32x64x4x32xf32>
 // CHECK-SAME:    %[[ARGY:[0-9a-z]+]]: memref<32x64x4x32xf32>
 func.func @saxpy4d(%x: !type4d, %y: !type4d, %alpha : f32) -> !type4d {
-  %c32 = arith.constant 32 : index  
-  %c64 = arith.constant 64 : index  
-  %c4 = arith.constant 4 : index  
+  %c32 = arith.constant 32 : index
+  %c64 = arith.constant 64 : index
+  %c4 = arith.constant 4 : index
 //      CHECK:   %[[C32:.*]] = arith.constant 32 : index
 //      CHECK:   %[[C64:.*]] = arith.constant 64 : index
 //      CHECK:   %[[C4:.*]] = arith.constant 4 : index
 //      CHECK:   %[[C1:.*]] = arith.constant 1 : index
 //      CHECK:   gpu.launch blocks(%{{.*}}, %{{.*}}, %{{.*}}) in (%{{.*}} = %[[C32]], %{{.*}} = %[[C64]], %{{.*}} = %[[C1]]) threads(%{{.*}}, %{{.*}}, %{{.*}}) in (%{{.*}} = %[[C32]], %{{.*}} = %[[C4]], %{{.*}} = %[[C1]])
 //      CHECK:   %[[BLKX:.*]] = gpu.block_id  x
 //      CHECK:   %[[BLKY:.*]] = gpu.block_id  y
 //      CHECK:   %[[TIDX:.*]] = gpu.thread_id  x
 //      CHECK:   %[[TIDY:.*]] = gpu.thread_id  y
 //      CHECK:   memref.load %[[ARGX]][%[[BLKX]], %[[BLKY]], %[[TIDY]], %[[TIDX]]]
 //      CHECK:   memref.load %[[ARGY]][%[[BLKX]], %[[BLKY]], %[[TIDY]], %[[TIDX]]]
   scf.foreach_thread (%i, %j) in (%c32, %c64) {
     scf.foreach_thread (%k, %l) in (%c4, %c32) {
-      %4 = memref.load %x[%i, %j, %k, %l] : !type4d        
+      %4 = memref.load %x[%i, %j, %k, %l] : !type4d
       %5 = memref.load %y[%i, %j, %k, %l] : !type4d
       %6 = math.fma %alpha, %4, %5 : f32
       memref.store %6, %y[%i, %j, %k, %l] : !type4d
     }  {thread_dim_mapping = [1, 0, 2]}
   }  {thread_dim_mapping = [0, 1, 2]}
   return %y : !type4d
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %funcop = transform.structured.match ops{["func.func"]} in %arg0
     %gpuLaunch = transform.gpu.map_foreach_to_blocks %funcop { generate_gpu_launch }
     transform.gpu.map_nested_foreach_to_threads %gpuLaunch { blockDim = [32, 4, 1] }
   }
 }
 
 // -----
%name = gpu.launch async[%stream] blocks(%arg3, %arg4, %arg5) in (%arg9 = %one, %arg10 = %one, %arg11 = %one)
      }  {thread_dim_mapping = [1, 0, 2]}
     gpu.terminator
   }
   return %y : !type
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0
     transform.gpu.map_nested_foreach_to_threads %funcop { blockDim = [12, 9, 1], syncAfterDistribute = false }
   }
 }

mlir/test/Dialect/Linalg/multisize-tiling-full.mlir

 // RUN: mlir-opt --test-transform-dialect-interpreter --canonicalize %s | FileCheck %s
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
   // This implements a 2D multisize tiling with target sizes [3, 10].
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
     %1:3 = transform.structured.multitile_sizes %0 { dimension = 0, target_size = 3}
     %t:3 = transform.structured.multitile_sizes %0 { dimension = 1, target_size = 10}
     %2:2 = transform.structured.split %0 after %1#2 { dimension = 0 }
     %3:2 = transform.structured.tile %2#0 [%1#0]
     %4:2 = transform.structured.tile %2#1 [%1#1]
-    %5 = merge_handles %3#0, %4#0
-    %tt:3 = replicate num(%5) %t#0, %t#1, %t#2
+    %5 = merge_handles %3#0, %4#0 : !pdl.operation
+    %tt:3 = replicate num(%5) %t#0, %t#1, %t#2 : !pdl.operation, !pdl.operation, !pdl.operation, !pdl.operation
     %6:2 = transform.structured.split %5 after %tt#2 { dimension = 1 }
     transform.structured.tile %6#0 [0, %tt#0]
     transform.structured.tile %6#1 [0, %tt#1]
   }
 }
 
 func.func private @elem(%arg0: f32, %arg1: index, %arg2: index) -> f32
 

mlir/test/Dialect/Linalg/promote.mlir

 //       CHECK:           memref<?x?xf32, strided<[?, 1], offset: ?>>
 //
 //   CHECK-NOT:         memref.dealloc %[[tmpA]] : memref<32xi8>
 //   CHECK-NOT:         memref.dealloc %[[tmpB]] : memref<48xi8>
 //   CHECK-NOT:         memref.dealloc %[[tmpC]] : memref<24xi8>
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
     ^bb0(%arg1: !pdl.operation):
       %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
       %1 = transform.structured.promote %0 { use_alloca }
   }
 }
 
 // -----
 

 //       CHECK:           memref<?x?xf64, strided<[?, 1], offset: ?>>
 //
 //       CHECK:         memref.dealloc %[[tmpA_f64]] : memref<64xi8>
 //       CHECK:         memref.dealloc %[[tmpB_f64]] : memref<96xi8>
 //       CHECK:         memref.dealloc %[[tmpC_f64]] : memref<48xi8>
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
     ^bb0(%arg1: !pdl.operation):
       %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
       %1 = transform.structured.promote %0
   }
 }
 
 
 // -----
^bb0(%arg9: f32, %arg10: f32, %arg11: f32):
     linalg.yield %16 : f32
   }
 
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
     ^bb0(%arg1: !pdl.operation):
       %0 = transform.structured.match interface{LinalgOp} in %arg1
       %1 = transform.structured.promote %0
   }
 }

mlir/test/Dialect/Linalg/promotion_options.mlir

 //      CHECK:       memref.copy %[[T17]], %[[T21]]
 //      CHECK:       linalg.matmul ins(%[[T19]], %[[T12]]{{.*}} outs(%[[T21]]
 //      CHECK:       memref.copy %[[T21]], %[[T17]]
 //      CHECK:       memref.dealloc %[[A0]]
 //      CHECK:       memref.dealloc %[[A1]]
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
     ^bb0(%arg1: !pdl.operation):
       %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
       %1, %loops:3 = transform.structured.tile %0 [16, 16, 16]
       %2 = transform.structured.promote %1 { operands_to_promote = [0, 2], force_full_tiles = [false, false] }
   }
 }

mlir/test/Dialect/Linalg/tile-to-foreach-thread.mlir

//  CHECK-SAME:   %[[C:[0-9a-z]+]]: tensor<?x?xf32>
   // CHECK-NEXT: } {thread_dim_mapping = [1, 0]}
     %0 = linalg.matmul ins(%A, %B : tensor<?x?xf32>, tensor<?x?xf32>)
                       outs(%C : tensor<?x?xf32>) -> (tensor<?x?xf32>)
     return %0 : tensor<?x?xf32>
   }
 
   transform.with_pdl_patterns {
   ^bb0(%arg0: !pdl.operation):
-    transform.sequence %arg0 failures(propagate) {
+    transform.sequence %arg0 : !pdl.operation failures(propagate) {
     ^bb1(%arg1: !pdl.operation):
       %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
       %1:2 = transform.structured.tile_to_foreach_thread_op %0 num_threads [10, 20] (mapped to dims [1, 0])
     }
   }
 }
 
 // -----
func.func @matmul_static(%A: tensor<100x200xf32>, %B: tensor<200x300xf32>, %C: tensor<100x300xf32>) -> tensor<100x300xf32> {
   // CHECK-NEXT:    tensor.parallel_insert_slice
   %0 = linalg.matmul ins(%A, %B : tensor<100x200xf32>, tensor<200x300xf32>)
                     outs(%C : tensor<100x300xf32>) -> (tensor<100x300xf32>)
   return %0 : tensor<100x300xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
     %1:2 = transform.structured.tile_to_foreach_thread_op %0 num_threads [10, 21]
   }
 }
 
 
 // -----
 
 // CHECK-DAG: #[[$map0:.+]] = affine_map<()[s0] -> (s0 ceildiv 10)>
 // CHECK-DAG: #[[$map1:.+]] = affine_map<()[s0] -> (s0 ceildiv 20)>
 // CHECK-DAG: #[[$map2:.+]] = affine_map<(d0)[s0] -> (d0 * -10 + s0, 10)>
 // CHECK-DAG: #[[$map4:.+]] = affine_map<(d0)[s0] -> (d0 * -20 + s0, 20)>
 // CHECK-DAG: #[[$map5:.+]] = affine_map<(d0) -> (d0 * 10)>
 // CHECK-DAG: #[[$map6:.+]] = affine_map<(d0) -> (d0 * 20)>
 
 // CHECK-LABEL: matmul_tile_size_dynamic(
 //  CHECK-SAME:   %[[A:[0-9a-z]+]]: tensor<?x?xf32>
 //  CHECK-SAME:   %[[B:[0-9a-z]+]]: tensor<?x?xf32>
 //  CHECK-SAME:   %[[C:[0-9a-z]+]]: tensor<?x?xf32>
-func.func @matmul_tile_size_dynamic(%A: tensor<?x?xf32>, %B: tensor<?x?xf32>, %C: tensor<?x?xf32>) -> tensor<?x?xf32> {  
+func.func @matmul_tile_size_dynamic(%A: tensor<?x?xf32>, %B: tensor<?x?xf32>, %C: tensor<?x?xf32>) -> tensor<?x?xf32> {
   //      CHECK: %[[M:.+]] = tensor.dim %[[A]], %c0 :
-  //      CHECK: %[[N:.+]] = tensor.dim %[[B]], %c1 : 
+  //      CHECK: %[[N:.+]] = tensor.dim %[[B]], %c1 :
   //      CHECK: %[[NT0:.+]] = affine.apply #map0()[%[[M]]]
   //      CHECK: %[[NT1:.+]] = affine.apply #map1()[%[[N]]]
   //      CHECK: scf.foreach_thread (%[[IV0:.+]], %[[IV1:.+]]) in (%[[NT0]], %[[NT1]]) shared_outs(%[[C_BLK:.*]] = %[[C]])
-  //      CHECK:   %[[TS0:.+]] = affine.min #[[$map2]](%[[IV0]])[%[[M]]]  
+  //      CHECK:   %[[TS0:.+]] = affine.min #[[$map2]](%[[IV0]])[%[[M]]]
   //      CHECK:   %[[TS1:.+]] = affine.min #[[$map4]](%[[IV1]])[%[[N]]]
   //      CHECK:   %[[LB0:.+]] = affine.apply #[[$map5]](%[[IV0]])
   //      CHECK:   %[[LB1:.+]] = affine.apply #[[$map6]](%[[IV1]])
   //      CHECK:   tensor.extract_slice %[[A]]
   //      CHECK:   tensor.extract_slice %[[B]]
   //      CHECK:   tensor.extract_slice %[[C_BLK]]
   //      CHECK:   linalg.matmul
   //      CHECK:   scf.foreach_thread.perform_concurrently
   // CHECK-NEXT:    tensor.parallel_insert_slice
   %0 = linalg.matmul ins(%A, %B : tensor<?x?xf32>, tensor<?x?xf32>)
                     outs(%C : tensor<?x?xf32>) -> (tensor<?x?xf32>)
   return %0 : tensor<?x?xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
-  ^bb1(%arg1: !pdl.operation):    
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
+  ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
     %1:2 = transform.structured.tile_to_foreach_thread_op %0 tile_sizes [10, 20]
   }
 }
 
 // -----
 
 // Tests that dimension 0 can eliminate affine.min/max, dimension 1 cannot.
 
 // CHECK-DAG: #[[$map0:.+]] = affine_map<(d0) -> (d0 * -21 + 300, 21)>
 // CHECK-DAG: #[[$map2:.+]] = affine_map<(d0) -> (d0 * 10)>
 // CHECK-DAG: #[[$map3:.+]] = affine_map<(d0) -> (d0 * 21)>
 
 // CHECK-LABEL: matmul_tile_size_static(
 //  CHECK-SAME:   %[[A:[0-9a-z]+]]: tensor
 //  CHECK-SAME:   %[[B:[0-9a-z]+]]: tensor
 //  CHECK-SAME:   %[[C:[0-9a-z]+]]: tensor
 func.func @matmul_tile_size_static(%A: tensor<100x200xf32>, %B: tensor<200x300xf32>, %C: tensor<100x300xf32>) -> tensor<100x300xf32> {
   //  CHECK-DAG: %[[c10:.+]] = arith.constant 10 :
   //  CHECK-DAG: %[[c15:.+]] = arith.constant 15 :
   //      CHECK: scf.foreach_thread (%[[IV0:.+]], %[[IV1:.+]]) in (%[[c10]], %[[c15]]) shared_outs(%[[C_BLK:.*]] = %[[C]])
-  //      CHECK:   %[[TS:.+]] = affine.min #[[$map0]](%[[IV1]])  
+  //      CHECK:   %[[TS:.+]] = affine.min #[[$map0]](%[[IV1]])
   //  CHECK-NOT:   affine.max
   //  CHECK-NOT:   affine.min
   //      CHECK:   %[[LB0:.+]] = affine.apply #[[$map2]](%[[IV0]])
   //      CHECK:   %[[LB1:.+]] = affine.apply #[[$map3]](%[[IV1]])
   //      CHECK:   %[[tA:.+]] = tensor.extract_slice %[[A]][%[[LB0]], 0] [10, 200] [1, 1] :
   //      CHECK:   %[[tB:.+]] = tensor.extract_slice %[[B]][0, %[[LB1]]] [200, %[[TS]]] [1, 1] :
   //      CHECK:   %[[tC:.+]] = tensor.extract_slice %[[C_BLK]][%[[LB0]], %[[LB1]]] [10, %[[TS]]] [1, 1] :
   //      CHECK:   linalg.matmul
   //      CHECK:   scf.foreach_thread.perform_concurrently
   // CHECK-NEXT:    tensor.parallel_insert_slice
   %0 = linalg.matmul ins(%A, %B : tensor<100x200xf32>, tensor<200x300xf32>)
                     outs(%C : tensor<100x300xf32>) -> (tensor<100x300xf32>)
   return %0 : tensor<100x300xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
     %1:2 = transform.structured.tile_to_foreach_thread_op %0 tile_sizes [10, 21]
   }
 }
 
 // -----
 
 module {
   func.func @extract_source(%A: tensor<4xf32>, %B: tensor<16xf32>) -> tensor<4xf32> {
     %B1 = tensor.extract_slice %B[10] [4] [1] : tensor<16xf32> to tensor<4xf32>
     %result = linalg.generic {indexing_maps = [
       affine_map<(d0) -> (d0)>,affine_map<(d0) -> (d0)>],
       iterator_types = ["parallel"]}
       ins(%A : tensor<4xf32>) outs(%B1 : tensor<4xf32>) {
       ^bb0(%arg3: f32, %arg4: f32):  // no predecessors
         %2 = arith.addf %arg3, %arg3 : f32
         linalg.yield %2 : f32
     } -> tensor<4xf32>
     return %result : tensor<4xf32>
   }
 
   transform.with_pdl_patterns {
   ^bb0(%arg0: !pdl.operation):
-    transform.sequence %arg0 failures(propagate) {
+    transform.sequence %arg0 : !pdl.operation failures(propagate) {
     ^bb1(%arg1: !pdl.operation):
       %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
       %1:2 = transform.structured.tile_to_foreach_thread_op %0 num_threads [2] (mapped to dims [0])
     }
   }
 }
 // CHECK-DAG: #[[$map0:.+]] = affine_map<(d0) -> (d0 * 2)>
 
func.func @matmul_tile_size_dynamic_dynamic(%A: tensor<?x?xf32>, %B: tensor<?x?xf32>, %C: tensor<?x?xf32>) -> tensor<?x?xf32> {
   %tile_size = "test.dummy"() : () -> (index)
   %0 = linalg.matmul ins(%A, %B : tensor<?x?xf32>, tensor<?x?xf32>)
                     outs(%C : tensor<?x?xf32>) -> (tensor<?x?xf32>)
   return %0 : tensor<?x?xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
     %sz = transform.structured.match ops{["test.dummy"]} in %arg1
     %1:2 = transform.structured.tile_to_foreach_thread_op %0 tile_sizes [%sz, 20]
   }
 }
 
 // -----
 
 // CHECK-DAG: #[[$map0:.+]] = affine_map<(d0) -> (d0 * -15 + 100, 15)>
 // CHECK-DAG: #[[$map1:.+]] = affine_map<(d0) -> (0, d0)>
 // CHECK-DAG: #[[$map2:.+]] = affine_map<(d0) -> (d0 * 15)>
 // CHECK-DAG: #[[$map3:.+]] = affine_map<(d0) -> (d0)>
 
 // CHECK-LABEL: tile_output_multi_1d_static(
 //  CHECK-SAME:   %[[IN1:[0-9a-z]+]]: tensor<100xf32>
 //  CHECK-SAME:   %[[IN2:[0-9a-z]+]]: tensor<100xf32>
 //  CHECK-SAME:   %[[ORGOUT1:[0-9a-z]+]]: tensor<100xf32>
 //  CHECK-SAME:   %[[ORGOUT2:[0-9a-z]+]]: tensor<100xf32>
-  func.func @tile_output_multi_1d_static(%IN1: tensor<100xf32>, %IN2: tensor<100xf32>, 
+  func.func @tile_output_multi_1d_static(%IN1: tensor<100xf32>, %IN2: tensor<100xf32>,
                                          %OUT1: tensor<100xf32>, %OUT2: tensor<100xf32>)
                                          -> (tensor<100xf32>, tensor<100xf32>) {
 //  CHECK-DAG: %[[c0:.+]] = arith.constant 7 :
 //      CHECK: scf.foreach_thread (%[[IV0:.+]]) in (%[[c0]]) shared_outs(%[[OUT1:[0-9a-z]+]] = %[[ORGOUT1]], %[[OUT2:[0-9a-z]+]] = %[[ORGOUT2]])
 //      CHECK:   %[[TSMIN:.+]] = affine.min #[[$map0]](%[[IV0]])
 //      CHECK:   %[[TS:.+]] = affine.max #[[$map1]](%[[TSMIN]])
 //  CHECK-NOT:   affine.min
 //  CHECK-NOT:   affine.max
// CHECK-NEXT:    tensor.parallel_insert_slice %[[RES1]]#1 into %[[OUT2]][%[[LB]]] [%[[TS]]] [1] :
     %res1, %res2 = linalg.generic
     {
       indexing_maps = [affine_map<(d0) -> (d0)>,
                        affine_map<(d0) -> (d0)>,
                        affine_map<(d0) -> (d0)>,
                        affine_map<(d0) -> (d0)>],
       iterator_types = ["parallel"]
     } ins(%IN1, %IN2 : tensor<100xf32>, tensor<100xf32>)
-      outs(%OUT1, %OUT2 : tensor<100xf32>, tensor<100xf32>) 
+      outs(%OUT1, %OUT2 : tensor<100xf32>, tensor<100xf32>)
     {
       ^bb0(%a1: f32, %a2: f32, %a3: f32, %a4: f32):
         %1 = arith.addf %a1, %a3 : f32
         %2 = arith.addf %a2, %a4 : f32
         linalg.yield %1, %2 : f32,f32
     } -> (tensor<100xf32>, tensor<100xf32>)
     return %res1, %res2 : tensor<100xf32>, tensor<100xf32>
   }
 
   transform.with_pdl_patterns {
     ^bb0(%arg0: !pdl.operation):
-    transform.sequence %arg0 failures(propagate) {
+    transform.sequence %arg0 : !pdl.operation failures(propagate) {
       ^bb1(%arg1: !pdl.operation):
       %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
       %foreach_thread, %tiled_generic = transform.structured.tile_to_foreach_thread_op %0 num_threads [7]
     }
   }
 
 // -----
 
 // CHECK-DAG: #[[$map0:.+]] = affine_map<(d0) -> (d0 * 75)>
 // CHECK-DAG: #[[$map1:.+]] = affine_map<(d0, d1) -> (d1)>
 // CHECK-DAG: #[[$map2:.+]] = affine_map<(d0, d1) -> (d1, d0)
 // CHECK-DAG: #[[$map3:.+]] = affine_map<(d0, d1) -> (d0)>
 // CHECK-DAG: #[[$map4:.+]] = affine_map<(d0, d1) -> (d0, d1)>
 
 // CHECK-LABEL: tile_output_multi_1d2d_static(
 //  CHECK-SAME:   %[[IN1:[0-9a-z]+]]: tensor<100xf32>
 //  CHECK-SAME:   %[[IN2:[0-9a-z]+]]: tensor<100x300xf32>
 //  CHECK-SAME:   %[[IN3:[0-9a-z]+]]: tensor<300xf32>
 //  CHECK-SAME:   %[[ORGOUT1:[0-9a-z]+]]: tensor<300x100xf32>
 //  CHECK-SAME:   %[[ORGOUT2:[0-9a-z]+]]: tensor<300xf32>
-  func.func @tile_output_multi_1d2d_static(%IN1: tensor<100xf32>, %IN2: tensor<100x300xf32>, %IN3: tensor<300xf32>,                     
+  func.func @tile_output_multi_1d2d_static(%IN1: tensor<100xf32>, %IN2: tensor<100x300xf32>, %IN3: tensor<300xf32>,
                      %OUT1: tensor<300x100xf32>, %OUT2: tensor<300xf32>)
                      -> (tensor<300x100xf32>, tensor<300xf32>) {
 //  CHECK-DAG: %[[c0:.+]] = arith.constant 4 :
 //      CHECK: scf.foreach_thread (%[[IV0:.+]]) in (%[[c0]]) shared_outs(%[[OUT1:[0-9a-z]+]] = %[[ORGOUT1]], %[[OUT2:[0-9a-z]+]] = %[[ORGOUT2]])
 //      CHECK:   %[[LB:.+]] = affine.apply #[[$map0]](%[[IV0]])
 //      CHECK:   %[[tIN1:.+]] = tensor.extract_slice %[[IN2]][0, %[[LB]]] [100, 75]
 //      CHECK:   %[[tIN2:.+]] = tensor.extract_slice %[[IN3]][%[[LB]]] [75]
 //      CHECK:   %[[tOUT1:.+]] = tensor.extract_slice %[[OUT1]][%[[LB]], 0] [75, 100]
 //      CHECK:   %[[tOUT2:.+]] = tensor.extract_slice %[[OUT2]][%[[LB]]] [75]
 //      CHECK:   %[[RES1:[0-9]+]]:[[RES2:[0-9]+]] = linalg.generic
 //      CHECK:   scf.foreach_thread.perform_concurrently
 // CHECK-NEXT:    tensor.parallel_insert_slice %[[RES1]]#0 into %[[OUT1]][%[[LB]], 0] [75, 100]
 // CHECK-NEXT:    tensor.parallel_insert_slice %[[RES1]]#1 into %[[OUT2]][%[[LB]]] [75]
     %res2, %res3 = linalg.generic {
-      indexing_maps = [affine_map<(d0,d1) -> (d1)>,   
+      indexing_maps = [affine_map<(d0,d1) -> (d1)>,
                        affine_map<(d0,d1) -> (d1,d0)>,
-                       affine_map<(d0,d1) -> (d0)>,   
+                       affine_map<(d0,d1) -> (d0)>,
                        affine_map<(d0,d1) -> (d0,d1)>,
-                       affine_map<(d0,d1) -> (d0)>    
+                       affine_map<(d0,d1) -> (d0)>
                        ],
       iterator_types = ["parallel", "parallel"]
     } ins(%IN1, %IN2, %IN3 : tensor<100xf32>, tensor<100x300xf32>, tensor<300xf32>)
       outs(%OUT1, %OUT2: tensor<300x100xf32>, tensor<300xf32>)  {
       ^bb0(%i1: f32, %i2: f32, %i3: f32, %o1: f32, %o2: f32):
         %1 = arith.addf %i1, %o1 : f32
         %2 = arith.addf %i2, %1 : f32
         %3 = arith.addf %i3, %2 : f32
         linalg.yield %3, %i3 : f32, f32
     } -> (tensor<300x100xf32>, tensor<300xf32>)
-    
     return %res2, %res3 : tensor<300x100xf32>, tensor<300xf32>
   }
 
   transform.with_pdl_patterns {
     ^bb0(%IN_MAT1: !pdl.operation):
-    transform.sequence %IN_MAT1 failures(propagate) {
+    transform.sequence %IN_MAT1 : !pdl.operation failures(propagate) {
       ^bb1(%IN_MAT2: !pdl.operation):
       %0 = transform.structured.match ops{["linalg.generic"]} in %IN_MAT2
       %foreach_thread, %tiled_generic = transform.structured.tile_to_foreach_thread_op %0 num_threads [4]
     }
   }
 
 
 

mlir/test/Dialect/Linalg/transform-op-decompose.mlir

%0 = linalg.depthwise_conv_2d_nhwc_hwc {dilations = dense<1> : vector<2xi64>, strides = dense<2> : vector<2xi64>}
          ins(%input, %filter: tensor<1x1x113x96xf32>, tensor<1x3x96xf32>)
          outs(%init: tensor<1x1x56x96xf32>) -> tensor<1x1x56x96xf32>
   // CHECK: %[[INSERTED]]
   return %0: tensor<1x1x56x96xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match interface{LinalgOp} in %arg1
     %1 = transform.structured.decompose %0
   }
 }

mlir/test/Dialect/Linalg/transform-op-fuse-into-containing.mlir

module {
 
   // Check no failure when nothing happens.
   func.func @dummy1() { return }
   func.func @dummy2() { return }
   func.func @dummy3() { return }
 
   transform.with_pdl_patterns {
   ^bb0(%arg0: !pdl.operation):
-    transform.sequence %arg0 failures(propagate) {
+    transform.sequence %arg0 : !pdl.operation failures(propagate) {
     ^bb1(%arg1: !pdl.operation):
       %0 = transform.structured.match ops{["linalg.fill"]} in %arg1
       %1 = transform.structured.match ops{["scf.foreach_thread"]} in %arg1
 
       // linalg.fill is tileable. The op is tiled and fused.
       transform.structured.fuse_into_containing_op %0 into %1
     }
   }
func.func @fuse_untileable_op(%arg0: index, %arg1: tensor<64xf32>, %arg2: tensor<64xf32>) -> tensor<64xf32> {
     }
     // CHECK: }
 
     func.return %2 : tensor<64xf32>
   }
 
   transform.with_pdl_patterns {
   ^bb0(%arg0: !pdl.operation):
-    transform.sequence %arg0 failures(propagate) {
+    transform.sequence %arg0 : !pdl.operation failures(propagate) {
     ^bb1(%arg1: !pdl.operation):
       %0 = transform.structured.match ops{["tensor.empty"]} in %arg1
       %1 = transform.structured.match ops{["scf.foreach_thread"]} in %arg1
 
       // tensor.empty is not tileable. The op is cloned and fused.
       transform.structured.fuse_into_containing_op %0 into %1
     }
   }

mlir/test/Dialect/Linalg/transform-op-fuse.mlir

%0 = linalg.elemwise_unary ins(%arg0 : tensor<?x?xf32>)
                              outs(%arg1: tensor<?x?xf32>) -> tensor<?x?xf32>
   %1 = linalg.elemwise_binary ins(%0, %arg0 : tensor<?x?xf32>, tensor<?x?xf32>)
                              outs(%arg1: tensor<?x?xf32>) -> tensor<?x?xf32>
   return %1 : tensor<?x?xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.elemwise_binary"]} in %arg1
     %1, %loops:2 = transform.structured.fuse %0 {tile_sizes = [32, 32], tile_interchange = [0, 1]}
   }
 }
 
 // -----
 
%0 = linalg.elemwise_unary ins(%arg0 : tensor<?x?xf32>)
                              outs(%arg1: tensor<?x?xf32>) -> tensor<?x?xf32>
   %1 = linalg.elemwise_binary ins(%0, %arg0 : tensor<?x?xf32>, tensor<?x?xf32>)
                              outs(%arg1: tensor<?x?xf32>) -> tensor<?x?xf32>
   return %1 : tensor<?x?xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.elemwise_binary"]} in %arg1
     %1, %loops:2 = transform.structured.fuse %0 {tile_sizes = [32, 32], tile_interchange = [0, 1]}
     transform.loop.peel %loops#0
   }
 }
 
 // -----
^bb0(%arg0: f32, %arg1: f32):
     %2 = arith.addf %arg0, %arg1 : f32
     linalg.yield %2 : f32
   } -> tensor<12x25xf32>
   func.return %0 : tensor<12x25xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
     %1, %loops:2 = transform.structured.fuse %0 {tile_sizes = [5, 0, 7], tile_interchange = [0, 2, 1]}
     %2, %loops_2 = transform.structured.tile %1 [0, 4]
   }
 }

mlir/test/Dialect/Linalg/transform-op-generalize.mlir

 // RUN: mlir-opt --test-transform-dialect-interpreter %s | FileCheck %s
 
 // CHECK-LABEL: func.func @generalize_unary
 func.func @generalize_unary(%arg0: tensor<?x?xf32>, %arg1: tensor<?x?xf32>) -> tensor<?x?xf32> {
 
   // CHECK-NOT:   linalg.elemwise_unary
   //     CHECK:   linalg.generic
   %0 = linalg.elemwise_unary ins(%arg0 : tensor<?x?xf32>)
                              outs(%arg1: tensor<?x?xf32>) -> tensor<?x?xf32>
   return %0 : tensor<?x?xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.elemwise_unary"]} in %arg1
     %1 = transform.structured.generalize %0
   }
 }

mlir/test/Dialect/Linalg/transform-op-interchange.mlir

^bb0(%arg2: f32, %arg3: f32):
     %1 = math.exp %arg2 : f32
     linalg.yield %1 : f32
   } -> tensor<?x?xf32>
   return %0 : tensor<?x?xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
     transform.structured.interchange %0 { iterator_interchange = [1, 0]}
   }
 }
 
 // -----
 
 func.func @interchange_matmul(%arg0: tensor<?x?xf32>, %arg1: tensor<?x?xf32>, %arg2: tensor<?x?xf32>) -> tensor<?x?xf32> {
   // expected-note @below {{when applied to this op}}
   %0 = linalg.matmul ins(%arg0, %arg1 : tensor<?x?xf32>, tensor<?x?xf32>) outs(%arg2 : tensor<?x?xf32>) -> tensor<?x?xf32>
   return %0 : tensor<?x?xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
     // expected-error @below {{transform applied to the wrong op kind}}
     transform.structured.interchange %0 { iterator_interchange = [1, 0]}
   }
 }

mlir/test/Dialect/Linalg/transform-op-match.mlir

 // RUN: mlir-opt %s --test-transform-dialect-interpreter -allow-unregistered-dialect --split-input-file --verify-diagnostics
 
 func.func @bar() {
   // expected-remark @below {{matched op name}}
   // expected-remark @below {{matched attr name}}
   %0 = arith.constant {my_attr} 0: i32
   // expected-remark @below {{matched op name}}
   %1 = arith.constant 1 : i32
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %match_name = transform.structured.match ops{["arith.constant"]} in %arg1
     transform.test_print_remark_at_operand %match_name, "matched op name"
     transform.test_consume_operand %match_name
 
     %match_attr = transform.structured.match ops{["arith.constant"]} attributes{my_attr} in %arg1
     transform.test_print_remark_at_operand %match_attr, "matched attr name"
     transform.test_consume_operand %match_attr
   }
 }
 
 // -----
 
 func.func @by_type() {
   %0 = arith.constant 0: i32
   // expected-remark @below {{matched op name}}
   %1 = arith.constant 1.0 : f32
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %match_name = transform.structured.match
       ops{["arith.constant"]} filter_result_type = f32 in %arg1
     transform.test_print_remark_at_operand %match_name, "matched op name"
     transform.test_consume_operand %match_name
   }
 }
 
func.func @match_complex_attribute(%arg0: tensor<12x128x32xf32>)
   ^bb0(%arg1: f32, %arg2: f32):
     linalg.yield %arg1 : f32
   } -> tensor<128x12x32xf32>
   return %1 : tensor<128x12x32xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %match_attr = transform.structured.match
         ops{["linalg.generic"]}
         attributes{iterator_types = ["parallel", "parallel", "parallel"]}
         in %arg1
     transform.test_print_remark_at_operand %match_attr, "matched complex attr"
     transform.test_consume_operand %match_attr
 
     %no_match = transform.structured.match
         attributes{iterator_types = ["parallel", "parallel", "reduction"]}
         in %arg1
   // expected-remark @below {{0}}
     transform.test_print_number_of_associated_payload_ir_ops %no_match
   }
 }

mlir/test/Dialect/Linalg/transform-op-multitile-sizes.mlir

 // RUN: mlir-opt %s --test-transform-dialect-interpreter --split-input-file | FileCheck %s
 
 // CHECK-DAG: #[[$MAP13:.+]] = affine_map<() -> (13)>
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
     ^bb0(%arg1: !pdl.operation):
       %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
       transform.structured.multitile_sizes %0 { target_size = 3, dimension = 0 }
   }
 }
 
 // CHECK-LABEL: @multitile_sizes_static
 func.func @multitile_sizes_static(
func.func @multitile_sizes_static(
 
   return %0 : tensor<13x42xf32>
 }
 
 // -----
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
     ^bb0(%arg1: !pdl.operation):
       %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
       transform.structured.multitile_sizes %0 { target_size = 3, divisor = 2, dimension = 0 }
   }
 }
 
 // CHECK: #[[$MAP_A:.+]] = affine_map<()[s0] -> ([[A_IMPL:s0 floordiv 2]])>
 // CHECK: #[[$MAP_T:.+]] = affine_map<() -> (2)>

mlir/test/Dialect/Linalg/transform-op-pad.mlir

func.func @static_sizes_output_divisible(%arg0: tensor<24x12xf32>,
   // CHECK-SAME:              outs(%[[T2]] : tensor<4x5xf32>)
   %4 = linalg.matmul ins(%1, %2 : tensor<4x?xf32>, tensor<?x5xf32>) outs(%3 : tensor<4x5xf32>) -> tensor<4x5xf32>
   %5 = tensor.insert_slice %4 into %arg2[%iv0, %iv1] [4, 5] [1, 1] : tensor<4x5xf32> into tensor<24x25xf32>
   func.return %5 : tensor<24x25xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
     %1 = transform.structured.pad %0 {padding_values=[0.0 : f32, 0.0 : f32, 0.0 : f32], padding_dimensions=[0, 1, 2], pack_paddings=[1, 1, 0]}
   }
 }
 
 // -----
 
 func.func @pad(%arg0: tensor<24x12xf32>,
                %arg1: tensor<12x25xf32>,
                %arg2: tensor<24x25xf32>) -> tensor<24x25xf32> {
   // expected-note @below {{when applied to this op}} 
   %0 = linalg.matmul ins(%arg0, %arg1 : tensor<24x12xf32>, tensor<12x25xf32>) outs(%arg2 : tensor<24x25xf32>) -> tensor<24x25xf32>
   func.return %0 : tensor<24x25xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
     // expected-error @below {{op expects a padding value of type 'f32', got 0 : i32}}
     %1 = transform.structured.pad %0 {padding_values=[0: i32, 0.0 : f32, 0.0 : f32], padding_dimensions=[0, 1, 2], pack_paddings=[1, 1, 0]}
   }
 }
 
 // -----
 
 func.func @pad(%arg0: tensor<24x12xf32>,
                %arg1: tensor<12x25xf32>,
                %arg2: tensor<24x25xf32>) -> tensor<24x25xf32> {
   // expected-note @below {{when applied to this op}}
   %0 = linalg.matmul ins(%arg0, %arg1 : tensor<24x12xf32>, tensor<12x25xf32>) outs(%arg2 : tensor<24x25xf32>) -> tensor<24x25xf32>
   func.return %0 : tensor<24x25xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
     // expected-error @below {{expects a padding that parses to 'f32', got "foo"}}
     %1 = transform.structured.pad %0 {padding_values=["foo", 0.0 : f32, 0.0 : f32], padding_dimensions=[0, 1, 2], pack_paddings=[1, 1, 0]}
   }
 }
 
 // -----
 
 func.func @pad(%arg0: tensor<24x12xf32>,
                %arg1: tensor<12x25xf32>,
                %arg2: tensor<24x25xf32>) -> tensor<24x25xf32> {
   // This is attached to an error that is silenceable and is not reported by this transform
   //   {{when applied to this op}}
   %0 = linalg.matmul ins(%arg0, %arg1 : tensor<24x12xf32>, tensor<12x25xf32>) outs(%arg2 : tensor<24x25xf32>) -> tensor<24x25xf32>
   func.return %0 : tensor<24x25xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(suppress) {
+  transform.sequence %arg0 : !pdl.operation failures(suppress) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
     // This error is silenceable and is not reported by this transform
     //   {{transform.structured.pad failed to apply}}
     %1 = transform.structured.pad %0 {padding_values=[0.0 : f32, 0.0 : f32, 0.0 : f32], padding_dimensions=[0, 1, 2], pack_paddings=[1, 1, 0]}
   }
 }

mlir/test/Dialect/Linalg/transform-op-scalarize.mlir

 // RUN: mlir-opt -test-transform-dialect-interpreter %s | FileCheck %s
 
 func.func @scalarize(%arg0: tensor<24x12xf32>,
                      %arg1: tensor<12x25xf32>,
                      %arg2: tensor<24x25xf32>) -> tensor<24x25xf32> {
   // The op is first tiled by 10 in the first dimension, which creates a
   // dynamic size, and then scalarized, which brings the dimension to static 1.
   // CHECK: linalg.matmul ins(%{{.*}}, %{{.*}} : tensor<1x12
   %0 = linalg.matmul ins(%arg0, %arg1 : tensor<24x12xf32>, tensor<12x25xf32>) outs(%arg2 : tensor<24x25xf32>) -> tensor<24x25xf32>
   func.return %0 : tensor<24x25xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
     %1, %loops = transform.structured.tile %0 [10, 0, 0]
     %2 = transform.structured.scalarize %1
   }
 }

mlir/test/Dialect/Linalg/transform-op-split-reduction-by-scaling.mlir

func.func @matmul_split(%A : tensor<?x256xf32>, %B: tensor<256x32xf32>, %C: tensor<?x32xf32>) -> tensor<?x32xf32> {
   // CHECK-SAME: outs(%{{[a-zA-Z0-9]*}} : tensor<?x32xf32>) {
   %0 = linalg.matmul ins(%A, %B: tensor<?x256xf32>, tensor<256x32xf32>)
                     outs(%C: tensor<?x32xf32>) -> tensor<?x32xf32>
   return %0: tensor<?x32xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
     %1:4 = transform.structured.split_reduction %0 
       { split_factor = 4, insert_split_dimension = 2, use_scaling_algorithm, use_alloc}
   }
 }

mlir/test/Dialect/Linalg/transform-op-split-reduction.mlir

func.func @matmul_split(%A : tensor<16x256xf32>, %B: tensor<256x32xf32>, %C: tensor<16x32xf32>) -> tensor<16x32xf32> {
   // CHECK-SAME: outs(%{{[a-zA-Z0-9_]*}} : tensor<16x32xf32>) {
   %0 = linalg.matmul ins(%A, %B: tensor<16x256xf32>, tensor<256x32xf32>)
                     outs(%C: tensor<16x32xf32>) -> tensor<16x32xf32>
   return %0: tensor<16x32xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
     %1:4 = transform.structured.split_reduction %0 { split_factor = 4, insert_split_dimension = 2}
   }
 }

mlir/test/Dialect/Linalg/transform-op-split.mlir

 // RUN: mlir-opt %s --test-transform-dialect-interpreter --split-input-file -verify-diagnostics | FileCheck %s
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
     %1:2 = transform.structured.split %0 after 42 { dimension = 0 }
   }
 }
 
 func.func private @elem(%arg0: f32, %arg1: index, %arg2: index) -> f32
 
func.func @one_d_static_overflow(%arg0: tensor<10xf32>, %arg1: tensor<10xf32>) -> tensor<10xf32> {
   } -> tensor<10xf32>
   return %0 : tensor<10xf32>
 }
 
 // -----
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
     %1 = transform.structured.match ops{["func.call"]} in %arg1
     transform.structured.split %0 after %1 { dimension = 0 }
   }
 }
 
 func.func private @get_size() -> index
func.func @dynamic(%arg0: tensor<100xf32>, %arg1: tensor<100xf32>) -> tensor<100xf32> {
   } -> tensor<100xf32>
   return %1 : tensor<100xf32>
 }
 
 // -----
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
     %1:2 = transform.structured.split %0 after 4 { dimension = 0}
     %2:2 = transform.structured.split %1#1 after 16 { dimension = 1 }
   }
 }
 
 func.func private @elem(%arg0: f32, %arg1: index, %arg2: index) -> f32
^bb1(%arg1: !pdl.operation):
   // expected-error @below {{expects either a dynamic or a static split point to be provided}}
   %0:2 = "transform.structured.split"(%arg1) { dimension = 1, static_split_point = -1 } : (!pdl.operation) -> (!pdl.operation, !pdl.operation)
 }
 
 // -----
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
     %1 = transform.structured.match ops{["func.call"]} in %arg1
     // expected-error @below {{expected dynamic split point handle to point to a single-result index-typed op}}
     transform.structured.split %0 after %1 { dimension = 0 }
   }
 }
 
func.func @dynamic(%arg0: tensor<100xf32>, %arg1: tensor<100xf32>) -> tensor<100xf32> {
   } -> tensor<100xf32>
   return %1 : tensor<100xf32>
 }
 
 // -----
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
     %1 = transform.structured.match ops{["func.call"]} in %arg1
     // expected-error @below {{expected the dynamic split point handle to point to as many operations (0) as the target handle (1)}}
     transform.structured.split %0 after %1 { dimension = 0 }
   }
 }
 

 ^bb0(%arg0: !pdl.operation):
   pdl.pattern @func_return : benefit(1) {
     %0 = pdl.operands
     %1 = pdl.types
     %2 = pdl.operation "func.return"(%0 : !pdl.range<value>) -> (%1 : !pdl.range<type>)
     pdl.rewrite %2 with "transform.dialect"
   }
 
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["func.return"]} in %arg1
     // expected-error @below {{only applies to structured ops}}
     transform.structured.split %0 after 16 { dimension = 1 }
   }
 }
 
 func.func @noop(%arg0: tensor<100xf32>, %arg1: tensor<100xf32>) -> tensor<100xf32> {
   // expected-note @below {{target op}}
   return %arg0 : tensor<100xf32>
 }
 
 // -----
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
   pdl.pattern @linalg_generic : benefit(1) {
     %0 = pdl.operands
     %1 = pdl.types
     %2 = pdl.operation "linalg.generic"(%0 : !pdl.range<value>) -> (%1 : !pdl.range<type>)
     pdl.rewrite %2 with "transform.dialect"
   }
 
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
     // expected-error @below {{dimension 1 does not exist in target op}}
     transform.structured.split %0 after 16 { dimension = 1 }
   }
 }
 
 func.func @one_d_static(%arg0: tensor<100xf32>, %arg1: tensor<100xf32>) -> tensor<100xf32> {

mlir/test/Dialect/Linalg/transform-op-tile.mlir

 // RUN: mlir-opt --test-transform-dialect-interpreter --split-input-file %s | FileCheck %s
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
     ^bb0(%arg1: !pdl.operation):
       %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
       %1, %loops:3 = transform.structured.tile %0 [4, 4, 4]
   }
 }
 
 // CHECK-LABEL: func @tile_linalg_matmul(
 // CHECK-SAME:    %[[TA:[0-9a-z]+]]: tensor<128x128xf32>

 //      CHECK: return %[[TD0]] : tensor<128x128xf32>
   return %0 : tensor<128x128xf32>
 }
 
 // -----
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
     ^bb0(%arg1: !pdl.operation):
       %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
       %1 = transform.structured.match ops{["func.call"]} in %arg1
       %2, %loops:3 = transform.structured.tile %0 [%1, %1, 4]
   }
 }
 
 func.func private @get_dynamic_tile_size() -> index

mlir/test/Dialect/Linalg/transform-op-vectorize.mlir

^bb0(%arg0: !pdl.operation):
   pdl.pattern @pdl_target : benefit(1) {
     %args = operands
     %results = types
     %0 = pdl.operation "linalg.matmul"(%args : !pdl.range<value>) -> (%results : !pdl.range<type>)
     // TODO: we don't want this, but it is the required terminator for pdl.pattern
     rewrite %0 with "transform.dialect"
   }
 
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
-    %1 = get_closest_isolated_parent %0
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
     %2 = transform.structured.vectorize %1
   }
 }
 
 // -----
 
 #map0 = affine_map<()[s0] -> (-s0 + 12, 7)>
 #map1 = affine_map<()[s0] -> (-s0 + 7)>
func.func @vectorize_keep_pad(
   // CHECK: vector.transfer_write %[[vR]], %[[C]]
   %8 = linalg.matmul ins(%5, %7 : tensor<4x7xf32>, tensor<7x5xf32>) outs(%3 : tensor<4x5xf32>) -> tensor<4x5xf32>
   %9 = tensor.insert_slice %8 into %arg2[%arg3, %arg4] [4, 5] [1, 1] : tensor<4x5xf32> into tensor<24x25xf32>
   return %9 : tensor<24x25xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
-    %1 = get_closest_isolated_parent %0
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
     %2 = transform.structured.vectorize %1
   }
 }
 
 // -----
 
 #map0 = affine_map<()[s0] -> (-s0 + 12, 7)>
 #map1 = affine_map<()[s0] -> (-s0 + 7)>
func.func @vectorize_pad(
   // CHECK: vector.transfer_write %[[vR]], %[[C]]
   %8 = linalg.matmul ins(%5, %7 : tensor<4x7xf32>, tensor<7x5xf32>) outs(%3 : tensor<4x5xf32>) -> tensor<4x5xf32>
   %9 = tensor.insert_slice %8 into %arg2[%arg3, %arg4] [4, 5] [1, 1] : tensor<4x5xf32> into tensor<24x25xf32>
   return %9 : tensor<24x25xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
-    %1 = get_closest_isolated_parent %0
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
     %2 = transform.structured.vectorize %1 {vectorize_padding}
   }
 }
 
 // -----
 
 func.func @vectorize(%arg0: tensor<24x12xf32>,
                      %arg1: tensor<12x25xf32>,
                      %arg2: tensor<24x25xf32>) -> tensor<24x25xf32> {
   // expected-note @below {{non-isolated target}}
   %0 = linalg.matmul ins(%arg0, %arg1 : tensor<24x12xf32>, tensor<12x25xf32>) outs(%arg2 : tensor<24x25xf32>) -> tensor<24x25xf32>
   func.return %0 : tensor<24x25xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
     // expected-error @below {{op requires isolated-from-above targets}}
     %2 = transform.structured.vectorize %0
   }
 }

mlir/test/Dialect/Linalg/transform-patterns.mlir

^bb(%a: f32, %b: f32, %c: f32):
       %e = arith.addf %c, %d: f32
       linalg.yield %e: f32
   }
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
     transform.structured.interchange %0 { iterator_interchange = [1, 2, 0]}
   }
 }
 
 // CHECK-LABEL:  func @permute_generic
 // CHECK:        linalg.generic {

mlir/test/Dialect/Linalg/transform-promotion.mlir

 // CHECK:               memref.copy %[[s1]], %[[l1]] : memref<?x?xf32, strided{{.*}}> to memref<?x?xf32, strided{{.*}}>
 // CHECK:               memref.copy %[[s2]], %[[l2]] : memref<?x?xf32, strided{{.*}}> to memref<?x?xf32, strided{{.*}}>
 // CHECK:               linalg.matmul
 // CHECK-SAME:                 ins(%[[v0]], %[[v1]] : memref<?x?xf32>, memref<?x?xf32>)
 // CHECK-SAME:                outs(%[[v2]] : memref<?x?xf32>)
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
     ^bb0(%arg1: !pdl.operation):
       %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
       %1 = transform.structured.promote %0 { operands_to_promote = [0, 1, 2], use_full_tiles_by_default }
   }
 }
 
 // -----
 

 // CHECK:         memref.copy %[[s0]], %[[l0]] : memref<?x?xf32, strided{{.*}}> to memref<?x?xf32, strided{{.*}}>
 // CHECK-NOT:     memref.copy
 // CHECK:         linalg.matmul
 // CHECK-SAME:           ins(%[[v0]], %[[s1]] : memref<?x?xf32>, memref<?x?xf32, strided<[?, ?], offset: ?>>)
 // CHECK-SAME:          outs(%[[s2]] : memref<?x?xf32, strided<[?, ?], offset: ?>>)
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
     ^bb0(%arg1: !pdl.operation):
       %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
       %1 = transform.structured.promote %0 { operands_to_promote = [0], use_full_tiles_by_default }
   }
 }
 
 // -----
 

 // CHECK:         %[[v0:.*]] = memref.view %[[a0]]{{.*}} : memref<32000000xi8> to memref<?x?xf32>
 // CHECK:         %[[l0:.*]] = memref.subview %[[v0]][0, 0] [%{{.*}}, %{{.*}}] [1, 1] : memref<?x?xf32> to memref<?x?xf32, strided<[?, 1], offset: ?>>
 // CHECK:         linalg.fill ins({{.*}} : f32) outs(%[[v0]] : memref<?x?xf32>)
 // CHECK:         memref.copy %[[s0]], %[[l0]] : memref<?x?xf32, strided{{.*}}> to memref<?x?xf32, strided{{.*}}>
 // CHECK:         linalg.fill ins(%[[cf]] : f32) outs(%[[v0]] : memref<?x?xf32>)
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
     ^bb0(%arg1: !pdl.operation):
       %0 = transform.structured.match ops{["linalg.fill"]} in %arg1
       %1 = transform.structured.promote %0 { operands_to_promote = [1], use_full_tile_buffers = [false, true], alignment = 32}
   }
 }
 
 // -----
 

 // CHECK:         %[[v0:.*]] = memref.view %[[a0]]{{.*}} : memref<64000000xi8> to memref<?x?xcomplex<f32>>
 // CHECK:         %[[l0:.*]] = memref.subview %[[v0]][0, 0] [%{{.*}}, %{{.*}}] [1, 1] : memref<?x?xcomplex<f32>> to memref<?x?xcomplex<f32>, strided<[?, 1], offset: ?>>
 // CHECK:         linalg.fill ins({{.*}} : complex<f32>) outs(%[[v0]] : memref<?x?xcomplex<f32>>)
 // CHECK:         memref.copy %[[s0]], %[[l0]] : memref<?x?xcomplex<f32>, strided{{.*}}> to memref<?x?xcomplex<f32>, strided{{.*}}>
 // CHECK:         linalg.fill ins(%[[cc]] : complex<f32>) outs(%[[v0]] : memref<?x?xcomplex<f32>>)
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
     ^bb0(%arg1: !pdl.operation):
       %0 = transform.structured.match ops{["linalg.fill"]} in %arg1
       %1 = transform.structured.promote %0 { operands_to_promote = [1], use_full_tile_buffers = [false, true], alignment = 32}
   }
 }

mlir/test/Dialect/Linalg/transform-tile-and-fuse.mlir

^bb0(%arg2: f32, %arg3: f32, %arg4: f32):
       %18 = arith.select %17, %cst, %16 : f32
       linalg.yield %18 : f32
     } -> tensor<?x?xf32>
     return %7 : tensor<?x?xf32>
   }
 
   transform.with_pdl_patterns {
   ^bb0(%arg0: !pdl.operation):
-    transform.sequence %arg0 failures(propagate) {
+    transform.sequence %arg0 : !pdl.operation failures(propagate) {
     ^bb1(%arg1: !pdl.operation):
       // Find the root and all producers.
       %root = transform.structured.match attributes{"__root__"} in %arg1
       %producers = transform.structured.match attributes{"__producer__"} in %arg1
 
       // Tile the root.
       %foreach_thread_op, %tiled_op = transform.structured.tile_to_foreach_thread_op %root num_threads [10, 20]
 
^bb0(%arg2: f32, %arg3: f32, %arg4: f32):
       %18 = arith.select %17, %cst, %16 : f32
       linalg.yield %18 : f32
     } -> tensor<?x?xf32>
     return %7 : tensor<?x?xf32>
   }
 
   transform.with_pdl_patterns {
   ^bb0(%arg0: !pdl.operation):
-    transform.sequence %arg0 failures(propagate) {
+    transform.sequence %arg0 : !pdl.operation failures(propagate) {
     ^bb1(%arg1: !pdl.operation):
       // Find the root and all producers.
       %root = transform.structured.match attributes{"__root__"} in %arg1
       %producers = transform.structured.match attributes{"__producer__"} in %arg1
       %reversed_producers = transform.test_reverse_payload_ops %producers
 
       // Tile the root.
       %foreach_thread_op, %tiled_op = transform.structured.tile_to_foreach_thread_op %root num_threads [10, 20]
 
       // Fuse all producers.
       transform.structured.fuse_into_containing_op %reversed_producers into %foreach_thread_op
     }
   }
 }

mlir/test/Dialect/Linalg/vectorization.mlir

 // RUN: mlir-opt %s -test-transform-dialect-interpreter -split-input-file | FileCheck %s
 
 // -----
 
 // CHECK-LABEL: contraction_dot
 func.func @contraction_dot(%A: memref<1584xf32>, %B: memref<1584xf32>, %C: memref<f32>) {
 
 // CHECK: arith.mulf %{{.*}}, %{{.*}} : vector<1584xf32>
 // CHECK: vector.multi_reduction <add>, %{{.*}}, {{.*}} [0] : vector<1584xf32> to f32
   linalg.dot ins(%A, %B: memref<1584xf32>, memref<1584xf32>)
             outs(%C: memref<f32>)
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.dot"]} in %arg1
-    %1 = get_closest_isolated_parent %0
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
     %2 = transform.structured.vectorize %1  { disable_multi_reduction_to_contract_patterns }
   }
 }
 
 // -----
 
 // CHECK-LABEL: contraction_matvec
 func.func @contraction_matvec(%A: memref<1584x1584xf32>, %B: memref<1584xf32>, %C: memref<1584xf32>) {
 
 // CHECK: arith.mulf %{{.*}}, %{{.*}} : vector<1584x1584xf32>
 // CHECK: vector.multi_reduction <add>, %{{.*}}, {{.*}} [1] : vector<1584x1584xf32> to vector<1584xf32>
   linalg.matvec ins(%A, %B: memref<1584x1584xf32>, memref<1584xf32>)
             outs(%C: memref<1584xf32>)
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.matvec"]} in %arg1
-    %1 = get_closest_isolated_parent %0
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
     %2 = transform.structured.vectorize %1  { disable_multi_reduction_to_contract_patterns }
   }
 }
 
 // -----
 
 // CHECK-LABEL: contraction_matmul
 func.func @contraction_matmul(%A: memref<1584x1584xf32>, %B: memref<1584x1584xf32>, %C: memref<1584x1584xf32>) {
 // CHECK: arith.mulf %{{.*}}, %{{.*}} : vector<1584x1584x1584xf32>
 // CHECK: vector.multi_reduction <add>, %{{.*}}, {{.*}} [2] : vector<1584x1584x1584xf32> to vector<1584x1584xf32>
   linalg.matmul ins(%A, %B: memref<1584x1584xf32>, memref<1584x1584xf32>)
             outs(%C: memref<1584x1584xf32>)
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
-    %1 = get_closest_isolated_parent %0
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
     %2 = transform.structured.vectorize %1  { disable_multi_reduction_to_contract_patterns }
   }
 }
 
 // -----
 
 // CHECK-LABEL: contraction_batch_matmul
 func.func @contraction_batch_matmul(%A: memref<1584x1584x1584xf32>, %B: memref<1584x1584x1584xf32>, %C: memref<1584x1584x1584xf32>) {
 // CHECK: arith.mulf %{{.*}}, %{{.*}} : vector<1584x1584x1584x1584xf32>
 // CHECK: vector.multi_reduction <add>, %{{.*}}, {{.*}} [3] : vector<1584x1584x1584x1584xf32> to vector<1584x1584x1584xf32>
   linalg.batch_matmul
     ins(%A, %B: memref<1584x1584x1584xf32>, memref<1584x1584x1584xf32>)
    outs(%C: memref<1584x1584x1584xf32>)
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.batch_matmul"]} in %arg1
-    %1 = get_closest_isolated_parent %0
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
     %2 = transform.structured.vectorize %1  { disable_multi_reduction_to_contract_patterns }
   }
 }
 
 // -----
 
 #matmul_trait = {
   args_in = 2,
^bb(%a: f32, %b: f32, %c: f32) :
       %e = arith.addf %c, %d: f32
       linalg.yield %e : f32
   }
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
-    %1 = get_closest_isolated_parent %0
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
     %2 = transform.structured.vectorize %1  { disable_multi_reduction_to_contract_patterns, disable_transfer_permutation_map_lowering_patterns }
   }
 }
 
 // -----
 
 #matmul_transpose_out_trait = {
   args_in = 2,
^bb(%a: f32, %b: f32, %c: f32) :
       %e = arith.addf %c, %d: f32
       linalg.yield %e : f32
   }
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
-    %1 = get_closest_isolated_parent %0
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
     %2 = transform.structured.vectorize %1  { disable_multi_reduction_to_contract_patterns, disable_transfer_permutation_map_lowering_patterns }
   }
 }
 
 // -----
 
 #map0 = affine_map<(d0, d1, d2) -> (d0, d1, d2)>
 #map1 = affine_map<(d0, d1, d2) -> (d1, d0, d2)>
func.func @generic_interchanged_transpose(%arg0: tensor<12x128x32xf32>) -> tensor<128x12x32xf32> {
   ^bb0(%arg1: f32, %arg2: f32):
     linalg.yield %arg1 : f32
   } -> tensor<128x12x32xf32>
   return %1 : tensor<128x12x32xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
-    %1 = get_closest_isolated_parent %0
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
     %2 = transform.structured.vectorize %1  { disable_multi_reduction_to_contract_patterns, disable_transfer_permutation_map_lowering_patterns }
   }
 }
 
 // -----
 
 #matmul_trait = {
   args_in = 2,
^bb(%a: i32, %b: i32, %c: i32) :
       %e = arith.addi %c, %d: i32
       linalg.yield %e : i32
   }
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
-    %1 = get_closest_isolated_parent %0
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
     %2 = transform.structured.vectorize %1  { disable_multi_reduction_to_contract_patterns, disable_transfer_permutation_map_lowering_patterns }
   }
 }
 
 // -----
 
 // CHECK-LABEL: func @vectorization_test_2
 func.func @vectorization_test_2(%A: memref<8x16xf32>, %B: memref<16x32xf32>,
                          %C: memref<8x32xf32>) {
   //       CHECK: arith.mulf %{{.*}}, %{{.*}} : vector<8x32x16xf32>
   //       CHECK: vector.multi_reduction <add>, %{{.*}}, {{.*}} [2] : vector<8x32x16xf32> to vector<8x32xf32>
   linalg.matmul
     ins(%A, %B: memref<8x16xf32>, memref<16x32xf32>)
    outs(%C: memref<8x32xf32>)
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
-    %1 = get_closest_isolated_parent %0
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
     %2 = transform.structured.vectorize %1  { disable_multi_reduction_to_contract_patterns }
   }
 }
 
 // -----
 
 // CHECK-LABEL: func @test_vectorize_scalar_input
 func.func @test_vectorize_scalar_input(%A : memref<8x16xf32>, %arg0 : f32) {
   //       CHECK: %[[V:.*]] = vector.broadcast {{.*}} : f32 to vector<8x16xf32>
   //       CHECK: vector.transfer_write %[[V]], {{.*}} : vector<8x16xf32>, memref<8x16xf32>
   linalg.generic {
     indexing_maps = [affine_map<(m, n) -> ()>, affine_map<(m, n) -> (m, n)>],
     iterator_types = ["parallel", "parallel"]}
    ins(%arg0 : f32)
   outs(%A: memref<8x16xf32>) {
     ^bb(%0: f32, %1: f32) :
       linalg.yield %0 : f32
   }
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
-    %1 = get_closest_isolated_parent %0
-    %2 = transform.structured.vectorize %1 
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
+    %2 = transform.structured.vectorize %1
   }
 }
 
 // -----
 
 // CHECK-LABEL: func @test_do_not_vectorize_unsupported_element_types
 func.func @test_do_not_vectorize_unsupported_element_types(%A : memref<8x16xcomplex<f32>>, %arg0 : complex<f32>) {
   // CHECK-NOT: vector.broadcast
   // CHECK-NOT: vector.transfer_write
   linalg.generic {
     indexing_maps = [affine_map<(m, n) -> ()>, affine_map<(m, n) -> (m, n)>],
     iterator_types = ["parallel", "parallel"]}
    ins(%arg0 : complex<f32>)
   outs(%A: memref<8x16xcomplex<f32>>) {
     ^bb(%0: complex<f32>, %1: complex<f32>) :
       linalg.yield %0 : complex<f32>
   }
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
-    %1 = get_closest_isolated_parent %0
-    %2 = transform.structured.vectorize %1 
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
+    %2 = transform.structured.vectorize %1
   }
 }
 
 // -----
 
 // CHECK-LABEL: func @test_vectorize_fill
 func.func @test_vectorize_fill(%A : memref<8x16xf32>, %arg0 : f32) {
   //       CHECK: %[[V:.*]] = vector.broadcast {{.*}} : f32 to vector<8x16xf32>
   //       CHECK: vector.transfer_write %[[V]], {{.*}} : vector<8x16xf32>, memref<8x16xf32>
   linalg.fill ins(%arg0 : f32) outs(%A : memref<8x16xf32>)
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.fill"]} in %arg1
-    %1 = get_closest_isolated_parent %0
-    %2 = transform.structured.vectorize %1 
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
+    %2 = transform.structured.vectorize %1
   }
 }
 
 // -----
 
 // CHECK-LABEL: func @test_vectorize_fill
 func.func @test_vectorize_fill_scalar(%A : memref<f32>, %arg0 : f32) {
   // CHECK-SAME: (%[[M:.*]]: memref<f32>, %[[val:.*]]: f32)
   //      CHECK:   %[[VEC:.*]] = vector.broadcast %[[val]] : f32 to vector<f32>
   //      CHECK:   vector.transfer_write %[[VEC]], %[[M]][] : vector<f32>, memref<f32>
   linalg.fill ins(%arg0 : f32) outs(%A : memref<f32>)
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.fill"]} in %arg1
-    %1 = get_closest_isolated_parent %0
-    %2 = transform.structured.vectorize %1 
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
+    %2 = transform.structured.vectorize %1
   }
 }
 
 // -----
 
 // CHECK-LABEL: func @test_vectorize_copy
 func.func @test_vectorize_copy(%A : memref<8x16xf32>, %B : memref<8x16xf32>) {
   //       CHECK: %[[V:.*]] = vector.transfer_read {{.*}} : memref<8x16xf32>, vector<8x16xf32>
   //       CHECK: vector.transfer_write %[[V]], {{.*}} : vector<8x16xf32>, memref<8x16xf32>
   memref.copy %A, %B :  memref<8x16xf32> to memref<8x16xf32>
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["memref.copy"]} in %arg1
-    %1 = get_closest_isolated_parent %0
-    %2 = transform.structured.vectorize %1 
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
+    %2 = transform.structured.vectorize %1
   }
 }
 
 // -----
 
 // CHECK-LABEL: func @test_vectorize_copy_scalar
 func.func @test_vectorize_copy_scalar(%A : memref<f32>, %B : memref<f32>) {
   //  CHECK-SAME: (%[[A:.*]]: memref<f32>, %[[B:.*]]: memref<f32>)
   //       CHECK:   %[[V:.*]] = vector.transfer_read %[[A]][]{{.*}} : memref<f32>, vector<f32>
   //       CHECK:   %[[val:.*]] = vector.extractelement %[[V]][] : vector<f32>
   //       CHECK:   %[[VV:.*]] = vector.broadcast %[[val]] : f32 to vector<f32>
   //       CHECK:   vector.transfer_write %[[VV]], %[[B]][] : vector<f32>, memref<f32>
   memref.copy %A, %B :  memref<f32> to memref<f32>
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["memref.copy"]} in %arg1
-    %1 = get_closest_isolated_parent %0
-    %2 = transform.structured.vectorize %1 
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
+    %2 = transform.structured.vectorize %1
   }
 }
 // -----
 
 // CHECK-LABEL: func @test_vectorize_trailing_index
   //  CHECK-SAME: (%[[ARG0:.*]]: memref<1x2x4x8xindex>)
 func.func @test_vectorize_trailing_index(%arg0: memref<1x2x4x8xindex>) {
   //   CHECK-DAG:   %[[CST0:.*]] = arith.constant dense<[0, 1, 2, 3, 4, 5, 6, 7]> : vector<8xindex>
//       CHECK:   vector.transfer_write %[[BCST]], %[[ARG0]][%[[C0]], %[[C0]], %[[C0]], %[[C0]]] {{.*}} : vector<1x2x4x8xindex>, memref<1x2x4x8xindex>
     %0 = linalg.index 3 : index
     linalg.yield %0 : index
   }
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
-    %1 = get_closest_isolated_parent %0
-    %2 = transform.structured.vectorize %1 
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
+    %2 = transform.structured.vectorize %1
   }
 }
 
 // -----
 
 // CHECK-LABEL: func @test_vectorize_inner_index
   //  CHECK-SAME: (%[[ARG0:.*]]: memref<1x2x4x8xindex>)
 func.func @test_vectorize_inner_index(%arg0: memref<1x2x4x8xindex>) {
//       CHECK:   vector.transfer_write %[[TRAN]], %[[ARG0]][%[[C0]], %[[C0]], %[[C0]], %[[C0]]] {{.*}} : vector<1x2x4x8xindex>, memref<1x2x4x8xindex>
     %0 = linalg.index 1 : index
     linalg.yield %0 : index
   }
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
-    %1 = get_closest_isolated_parent %0
-    %2 = transform.structured.vectorize %1 
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
+    %2 = transform.structured.vectorize %1
   }
 }
 
 // -----
 
 // CHECK-LABEL: func @generic_vectorize
   //  CHECK-SAME: (%[[ARG0:.*]]: memref<4x256xf32>, %[[ARG1:.*]]: memref<4x256xf32>,
   //  CHECK-SAME:  %[[ARG2:.*]]: memref<256xf32>, %[[ARG3:.*]]: f32)
//       CHECK:   vector.transfer_write %[[TAN]], %[[ARG0]][%[[C0]], %[[C0]]] {{.*}} : vector<4x256xf32>, memref<4x256xf32>
     linalg.yield %6, %8, %c1_f32, %9, %10, %11, %12, %13, %14, %15 : f32, f32,
       f32, f32, f32, f32, f32, f32, f32, f32
   }
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
-    %1 = get_closest_isolated_parent %0
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
     %2 = transform.structured.vectorize %1  { disable_transfer_permutation_map_lowering_patterns }
   }
 }
 
 // -----
 
 // CHECK-LABEL: func @generic_vectorize_tensor
 //  CHECK-SAME: (%[[ARG0:.*]]: tensor<4x256xf32>, %[[ARG1:.*]]: tensor<4x256xf32>,
func.func @generic_vectorize_tensor(%arg0: tensor<4x256xf32>,
   return %r#0, %r#1, %r#2, %r#3, %r#4, %r#5, %r#6, %r#7, %r#8, %r#9:
     tensor<4x256xf32>, tensor<4x256xf32>, tensor<4x256xf32>,
     tensor<4x256xf32>, tensor<4x256xf32>, tensor<4x256xf32>, tensor<4x256xf32>,
     tensor<4x256xf32>, tensor<4x256xf32>, tensor<4x256xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
-    %1 = get_closest_isolated_parent %0
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
     %2 = transform.structured.vectorize %1 { disable_transfer_permutation_map_lowering_patterns }
   }
 }
 
 // -----
 
 // CHECK-DAG: #[[$MAP0:.*]] = affine_map<(d0, d1) -> (d0, 0, 0, d1)>
 // CHECK-DAG: #[[$MAP1:.*]] = affine_map<(d0) -> (d0, 0, 0, 0)>
^bb0(%arg0: f32, %arg1: f32, %arg2: f32, %arg3: f32, %arg4: f32):
     %b = arith.addf %arg3, %a : f32
     linalg.yield %b : f32
   }
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
-    %1 = get_closest_isolated_parent %0
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
     %2 = transform.structured.vectorize %1  { disable_transfer_permutation_map_lowering_patterns }
   }
 }
 
 // -----
 
 // Test different input maps.
 #matmul_trait = {
^bb(%a: f32, %b: f32, %c: f32, %d: f32) :
       %f = arith.addf %e, %c: f32
       linalg.yield %f : f32
   }
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
-    %1 = get_closest_isolated_parent %0
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
     %2 = transform.structured.vectorize %1  { disable_transfer_permutation_map_lowering_patterns }
   }
 }
 
 // -----
 
 // CHECK-LABEL: func @matmul_tensors
 //  CHECK-SAME: (%[[ARG0:.*]]: tensor<8x4xf32>, %[[ARG1:.*]]: tensor<4x12xf32>,
%0 = linalg.matmul  ins(%arg0, %arg1: tensor<8x4xf32>, tensor<4x12xf32>)
                      outs(%arg2: tensor<8x12xf32>)
     -> tensor<8x12xf32>
   //       CHECK:   return %[[W]] : tensor<8x12xf32>
   return %0 : tensor<8x12xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
-    %1 = get_closest_isolated_parent %0
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
     %2 = transform.structured.vectorize %1  { disable_multi_reduction_to_contract_patterns, disable_transfer_permutation_map_lowering_patterns }
   }
 }
 
 // -----
 
 // CHECK-LABEL: func @pad_static(
 //  CHECK-SAME:                  %[[ARG0:.*]]: tensor<2x?x2xf32>, %[[PAD:.*]]: f32
^bb0(%arg1: index, %arg2: index, %arg3: index):
       tensor.yield %pad_value : f32
     } : tensor<2x?x2xf32> to tensor<2x3x4xf32>
   return %0 : tensor<2x3x4xf32>
 }
 
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["tensor.pad"]} in %arg1
-    %1 = get_closest_isolated_parent %0
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
     %2 = transform.structured.vectorize %1 { vectorize_padding }
   }
 }
 
 // -----
 
 // CHECK-LABEL: func @pad_static_source(
 //  CHECK-SAME:                  %[[ARG0:.*]]: tensor<2x5x2xf32>, %[[PAD:.*]]: f32
^bb0(%arg1: index, %arg2: index, %arg3: index):
       tensor.yield %pad_value : f32
     } : tensor<2x5x2xf32> to tensor<2x6x4xf32>
   return %0 : tensor<2x6x4xf32>
 }
 
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["tensor.pad"]} in %arg1
-    %1 = get_closest_isolated_parent %0
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
     %2 = transform.structured.vectorize %1  { vectorize_padding }
   }
 }
 
 
 // -----
 
 // CHECK-LABEL: func @pad_static_dynamic(
^bb0(%arg1: index, %arg2: index, %arg3: index, %arg4: index):
       tensor.yield %pad_value : f32
     } : tensor<1x2x2x?xf32> to tensor<6x?x?x?xf32>
   return %0 : tensor<6x?x?x?xf32>
 }
 
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["tensor.pad"]} in %arg1
-    %1 = get_closest_isolated_parent %0
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
     %2 = transform.structured.vectorize %1  { vectorize_padding }
   }
 }
 
 
 // -----
 
 // CHECK-LABEL: func @pad_and_transfer_read
func.func @pad_and_transfer_read(%arg0: tensor<5x6xf32>) -> vector<7x9xf32> {
   %1 = vector.transfer_read %0[%c0, %c0], %c6
       : tensor<10x13xf32>, vector<7x9xf32>
   return %1 : vector<7x9xf32>
 }
 
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["tensor.pad"]} in %arg1
-    %1 = get_closest_isolated_parent %0
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
     %2 = transform.structured.vectorize %1 { vectorize_padding }
   }
 }
 
 // -----
 
 func.func private @make_vector() -> vector<7x9xf32>
 
func.func @pad_and_transfer_write_static(
   %2 = vector.transfer_write %1, %0[%c0, %c0]
       : vector<7x9xf32>, tensor<10x13xf32>
   %3 = tensor.extract_slice %2[0, 0] [5, 6] [1, 1] : tensor<10x13xf32> to tensor<5x6xf32>
   return %3 : tensor<5x6xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %3 = transform.structured.match ops{["tensor.pad"]} in %arg1
-    %4 = get_closest_isolated_parent %3
+    %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
     %5 = transform.structured.vectorize %4  { vectorize_padding }
   }
 }
 
 
 // -----
 
 func.func private @make_vector() -> vector<7x9xf32>
func.func @pad_and_transfer_write_dynamic_static(
   %2 = vector.transfer_write %1, %0[%c0, %c0]
       : vector<7x9xf32>, tensor<?x13xf32>
   %3 = tensor.extract_slice %2[0, 0] [%size, 6] [1, 1] : tensor<?x13xf32> to tensor<?x6xf32>
   return %3 : tensor<?x6xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %3 = transform.structured.match ops{["tensor.pad"]} in %arg1
-    %4 = get_closest_isolated_parent %3
+    %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
     %5 = transform.structured.vectorize %4 { vectorize_padding }
   }
 }
 
 
 // -----
 
 func.func private @make_vector() -> tensor<12x13xf32>
func.func @pad_and_insert_slice_source(
   } : tensor<5x6xf32> to tensor<7x9xf32>
   %1 = call @make_vector() : () -> tensor<12x13xf32>
   %r = tensor.insert_slice %0 into %1[0, 0][7, 9][1, 1] : tensor<7x9xf32> into tensor<12x13xf32>
   return %r : tensor<12x13xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %3 = transform.structured.match ops{["tensor.pad"]} in %arg1
-    %4 = get_closest_isolated_parent %3
+    %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
     %5 = transform.structured.vectorize %4  { vectorize_padding }
   }
 }
 
 
 // -----
 
 func.func private @make_vector() -> tensor<12x13xf32>
func.func @pad_and_insert_slice_dest(
   } : tensor<1x5x6xf32> to tensor<1x12x13xf32>
   %1 = call @make_vector() : () -> tensor<12x13xf32>
   %r = tensor.insert_slice %1 into %0[0, 0, 0][1, 12, 13][1, 1, 1] : tensor<12x13xf32> into tensor<1x12x13xf32>
   return %r : tensor<1x12x13xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %3 = transform.structured.match ops{["tensor.pad"]} in %arg1
-    %4 = get_closest_isolated_parent %3
-    %5 = transform.structured.vectorize %4 
+    %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
+    %5 = transform.structured.vectorize %4
   }
 }
 
 // -----
 
 // CHECK-LABEL: func @pad_tensor_non_const_pad_value
 //  CHECK-SAME:     %[[ARG0:.*]]: tensor<5x6xf32>
 //   CHECK-NOT:   tensor.pad
^bb0(%arg1: index, %arg2: index):
       tensor.yield %m : f32
   } : tensor<5x6xf32> to tensor<12x13xf32>
   return %0 : tensor<12x13xf32>
 }
 
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %3 = transform.structured.match ops{["tensor.pad"]} in %arg1
-    %4 = get_closest_isolated_parent %3
+    %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
     %5 = transform.structured.vectorize %4  { vectorize_padding }
   }
 }
 
 // -----
 
 // CHECK-LABEL: func @sum_exp
 func.func @sum_exp(%input: tensor<4x16x8xf32>, %output: tensor<4x16xf32>)
^bb0(%arg0: f32, %arg1: f32):
       linalg.yield %2 : f32
     } -> tensor<4x16xf32>
   return %0 : tensor<4x16xf32>
 }
 
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %3 = transform.structured.match ops{["linalg.generic"]} in %arg1
-    %4 = get_closest_isolated_parent %3
-    %5 = transform.structured.vectorize %4 
+    %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
+    %5 = transform.structured.vectorize %4
   }
 }
 
 // -----
 
 // CHECK-DAG: #[[$M1:.*]] =  affine_map<(d0, d1) -> (d1, d0, 0, 0)>
 // CHECK-DAG: #[[$M2:.*]] =  affine_map<(d0, d1) -> (0, 0, d1, d0)>
 // CHECK-DAG: #[[$M3:.*]] =  affine_map<(d0, d1) -> (d1, d0)>
^bb0(%arg0: f32, %arg1: f32, %arg2: f32):
       linalg.yield %4 : f32
     } -> tensor<5x2xf32>
   return %0 : tensor<5x2xf32>
 }
 
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %3 = transform.structured.match ops{["linalg.generic"]} in %arg1
-    %4 = get_closest_isolated_parent %3
+    %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
     %5 = transform.structured.vectorize %4  { disable_multi_reduction_to_contract_patterns, disable_transfer_permutation_map_lowering_patterns }
   }
 }
 
 // -----
 
 // CHECK-LABEL:   func @red_max_2d(
 func.func @red_max_2d(%arg0: tensor<4x4xf32>) -> tensor<4xf32> {
^bb0(%in0: f32, %out0: f32):
     linalg.yield %max : f32
   } -> tensor<4xf32>
   return %red : tensor<4xf32>
 }
 
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %3 = transform.structured.match ops{["linalg.generic"]} in %arg1
-    %4 = get_closest_isolated_parent %3
+    %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
     %5 = transform.structured.vectorize %4 { vectorize_padding }
   }
 }
 
 // -----
 
 // CHECK-LABEL:   func @red_min_2d(
 func.func @red_min_2d(%arg0: tensor<4x4xf32>) -> tensor<4xf32> {
^bb0(%in0: f32, %out0: f32):
     linalg.yield %min : f32
   } -> tensor<4xf32>
   return %red : tensor<4xf32>
 }
 
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %3 = transform.structured.match ops{["linalg.generic"]} in %arg1
-    %4 = get_closest_isolated_parent %3
-    %5 = transform.structured.vectorize %4 
+    %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
+    %5 = transform.structured.vectorize %4
   }
 }
 
 // -----
 
 // CHECK-LABEL:   func @red_mul_2d(
 func.func @red_mul_2d(%arg0: tensor<4x4xf32>) -> tensor<4xf32> {
   // CHECK: tensor.empty() : tensor<4xf32>
^bb0(%in0: f32, %out0: f32):
     linalg.yield %mul : f32
   } -> tensor<4xf32>
   return %red : tensor<4xf32>
 }
 
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %3 = transform.structured.match ops{["linalg.generic"]} in %arg1
-    %4 = get_closest_isolated_parent %3
-    %5 = transform.structured.vectorize %4 
+    %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
+    %5 = transform.structured.vectorize %4
   }
 }
 
 // -----
 
 // CHECK-LABEL:   func @red_or_2d(
 func.func @red_or_2d(%arg0: tensor<4x4xi1>) -> tensor<4xi1> {
   // CHECK: tensor.empty() : tensor<4xi1>
^bb0(%in0: i1, %out0: i1):
     linalg.yield %or : i1
   } -> tensor<4xi1>
   return %red : tensor<4xi1>
 }
 
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %3 = transform.structured.match ops{["linalg.generic"]} in %arg1
-    %4 = get_closest_isolated_parent %3
-    %5 = transform.structured.vectorize %4 
+    %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
+    %5 = transform.structured.vectorize %4
   }
 }
 
 // -----
 
 // CHECK-LABEL:   func @red_and_2d(
 func.func @red_and_2d(%arg0: tensor<4x4xi1>) -> tensor<4xi1> {
   // CHECK: tensor.empty() : tensor<4xi1>
^bb0(%in0: i1, %out0: i1):
     linalg.yield %and : i1
   } -> tensor<4xi1>
   return %red : tensor<4xi1>
 }
 
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %3 = transform.structured.match ops{["linalg.generic"]} in %arg1
-    %4 = get_closest_isolated_parent %3
-    %5 = transform.structured.vectorize %4 
+    %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
+    %5 = transform.structured.vectorize %4
   }
 }
 
 // -----
 
 // CHECK-LABEL:   func @red_xor_2d(
 func.func @red_xor_2d(%arg0: tensor<4x4xi1>) -> tensor<4xi1> {
   // CHECK: tensor.empty() : tensor<4xi1>
^bb0(%in0: i1, %out0: i1):
     linalg.yield %xor : i1
   } -> tensor<4xi1>
   return %red : tensor<4xi1>
 }
 
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %3 = transform.structured.match ops{["linalg.generic"]} in %arg1
-    %4 = get_closest_isolated_parent %3
-    %5 = transform.structured.vectorize %4 
+    %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
+    %5 = transform.structured.vectorize %4
   }
 }
 
 // -----
 
 // CHECK-DAG: #[[$M5:.*]] = affine_map<(d0, d1) -> (d0, 0)>
 
 // CHECK-LABEL:   func @explicit_broadcast(
^bb0(%arg7: f32, %arg8: f32, %arg9: f32):
       linalg.yield %40 : f32
     } -> tensor<4x4xf32>
   return %red : tensor<4x4xf32>
 }
 
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %3 = transform.structured.match ops{["linalg.generic"]} in %arg1
-    %4 = get_closest_isolated_parent %3
-    %5 = transform.structured.vectorize %4 
+    %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
+    %5 = transform.structured.vectorize %4
   }
 }
 
 // -----
 
 // CHECK-DAG: #[[$M6:.*]] = affine_map<(d0, d1) -> (d0, 0)>
 
 // CHECK-LABEL:   func @fused_broadcast_red_2d
^bb0(%arg7: f32, %arg8: f32, %arg9: f32):
       linalg.yield %42 : f32
     } -> tensor<4xf32>
   return %red : tensor<4xf32>
 }
 
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.fill"]} in %arg1
-    %1 = get_closest_isolated_parent %0
-    %2 = transform.structured.vectorize %1 
-    
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
+    %2 = transform.structured.vectorize %1
     %3 = transform.structured.match ops{["linalg.generic"]} in %arg1
-    %4 = get_closest_isolated_parent %3
-    %5 = transform.structured.vectorize %4 
+    %4 = get_closest_isolated_parent %3 : (!pdl.operation) -> !pdl.operation
+    %5 = transform.structured.vectorize %4
   }
 }
 
 // -----
 
 //  CHECK-LABEL: func @reduce_1d(
 //   CHECK-SAME:   %[[A:.*]]: tensor<32xf32>
 func.func @reduce_1d(%arg0: tensor<32xf32>) -> tensor<f32> {
^bb0(%a: f32, %b: f32):
       linalg.yield %3 : f32
     } -> tensor<f32>
 
   return %2 : tensor<f32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
-    %1 = get_closest_isolated_parent %0
-    %2 = transform.structured.vectorize %1 
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
+    %2 = transform.structured.vectorize %1
   }
 }
 
 
 // -----
 
 // This test checks that vectorization does not occur when an input indexing map
-// is not a projected permutation. In the future, this can be converted to a 
+// is not a projected permutation. In the future, this can be converted to a
 // positive test when support is added.
 
 // CHECK-LABEL:   func @not_projected_permutation
 func.func @not_projected_permutation(%arg0: tensor<8x8xf32>) -> tensor<6x6x3x3xf32> {
   %c0 = arith.constant 0.0 : f32
   %init = tensor.empty() : tensor<6x6x3x3xf32>
   %fill = linalg.fill ins(%c0 : f32) outs(%init : tensor<6x6x3x3xf32>) -> tensor<6x6x3x3xf32>
   // CHECK: linalg.generic
   %result = linalg.generic {indexing_maps = [affine_map<(d0, d1, d2, d3) -> (d0 + d2, d1 + d3)>,
                                              affine_map<(d0, d1, d2, d3) -> (d0, d1, d2, d3)>],
    iterator_types = ["parallel", "parallel", "parallel", "parallel"]}
    ins(%arg0 : tensor<8x8xf32>)
    outs(%fill : tensor<6x6x3x3xf32>) {
     ^bb0(%arg7: f32, %arg9: f32):
       linalg.yield %arg7 : f32
     } -> tensor<6x6x3x3xf32>
   return %result : tensor<6x6x3x3xf32>
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
-    %1 = get_closest_isolated_parent %0
-    %2 = transform.structured.vectorize %1 
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
+    %2 = transform.structured.vectorize %1
   }
 }
 
 // -----
 
 // Check vectorization can handle cases where outputs are a mix of reduced and non-reduced values.
 func.func @mixed_parallel_reduced_results(%arg0 : tensor<2x4x8xf32>,
     %arg1 : tensor<2x4xf32>, %arg2 : tensor<2x4x8xf32>, %arg3 : tensor<2x4xf32>) ->

 //   CHECK-DAG:   %[[V2:.+]] = vector.transfer_read %[[ARG3]]
 //   CHECK-DAG:   %[[MUL:.+]] = arith.mulf %[[V0]], %[[V1]]
 //   CHECK-DAG:   %[[ADD:.+]] = vector.multi_reduction <add>, %[[MUL]], %[[V2]]
 //   CHECK-DAG:   vector.transfer_write %[[MUL]], %[[ARG2]]
 //   CHECK-DAG:   vector.transfer_write %[[ADD]], %[[ARG3]]
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["linalg.generic"]} in %arg1
-    %1 = get_closest_isolated_parent %0
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
     %2 = transform.structured.vectorize %1  { disable_multi_reduction_to_contract_patterns, disable_transfer_permutation_map_lowering_patterns }
   }
 }

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

scf.for %j = %arg0 to %arg1 step %arg2 {
       }
     }
   }
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["arith.addi"]} in %arg1
     // CHECK: = transform.loop.get_parent_for
     %1 = transform.loop.get_parent_for %0
     %2 = transform.loop.get_parent_for %0 { num_loops = 2 }
     %3 = transform.loop.get_parent_for %0 { num_loops = 3 }
     transform.test_print_remark_at_operand %1, "third loop"
     transform.test_print_remark_at_operand %2, "second loop"
     transform.test_print_remark_at_operand %3, "first loop"
   }
 }
 
 // -----
 
 func.func @get_parent_for_op_no_loop(%arg0: index, %arg1: index) {
   // expected-note @below {{target op}}
-  arith.addi %arg0, %arg1 : index  
+  arith.addi %arg0, %arg1 : index
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["arith.addi"]} in %arg1
     // expected-error @below {{could not find an 'scf.for' parent}}
     %1 = transform.loop.get_parent_for %0
   }
 }
 
 // -----
func.func @loop_outline_op(%arg0: index, %arg1: index, %arg2: index) {
   scf.for %j = %arg0 to %arg1 step %arg2 {
     arith.addi %j, %j : index
   }
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["arith.addi"]} in %arg1
     %1 = transform.loop.get_parent_for %0
     // CHECK: = transform.loop.outline %{{.*}}
     transform.loop.outline %1 {func_name = "foo"}
   }
 }
 
^bb0:
     func.call @body() : () -> ()
     scf.yield
   }
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["scf.while"]} in %arg1
     // expected-error @below {{failed to outline}}
     transform.loop.outline %0 {func_name = "foo"}
   }
 }
 
 // -----
func.func @loop_peel_op() {
   scf.for %i = %0 to %1 step %2 {
     arith.addi %i, %i : index
   }
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["arith.addi"]} in %arg1
     %1 = transform.loop.get_parent_for %0
     transform.loop.peel %1
   }
 }
 
 // -----
scf.for %i0 = %c0 to %c4 step %c1 {
     %A1_elem = arith.addf %A_elem, %cf : f32
     memref.store %A1_elem, %result[%i0] : memref<?xf32>
   }
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["arith.addf"]} in %arg1
     %1 = transform.loop.get_parent_for %0
     %2 = transform.loop.pipeline %1
     // Verify that the returned handle is usable.
     transform.test_print_remark_at_operand %2, "transformed"
   }
 }
scf.for %i = %c0 to %c42 step %c5 {
     // CHECK-COUNT-4: arith.addi %[[I]]
     arith.addi %i, %i : index
   }
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %0 = transform.structured.match ops{["arith.addi"]} in %arg1
     %1 = transform.loop.get_parent_for %0
     transform.loop.unroll %1 { factor = 4 }
   }
 }
 

mlir/test/Dialect/Transform/check-use-after-free.mlir

func.func @use_after_free_branching_control_flow() {
   ^bb1:
     // expected-note @below {{freed here}}
     transform.test_consume_operand_if_matches_param_or_fail %0[42]
     "transform.test_branching_transform_op_terminator"()[^bb3] : () -> ()
   ^bb2:
     "transform.test_branching_transform_op_terminator"()[^bb3] : () -> ()
   ^bb3:
     // expected-warning @below {{operand #0 may be used after free}}
-    transform.sequence %0 failures(propagate) {
+    transform.sequence %0 : !pdl.operation failures(propagate) {
     ^bb0(%arg0: !pdl.operation):
     }
     "transform.test_branching_transform_op_terminator"() : () -> ()
   }
   return
 }
 
 // -----
^bb1:
     transform.test_consume_operand_if_matches_param_or_fail %0[42]
     "transform.test_branching_transform_op_terminator"()[^bb3] : () -> ()
   ^bb2:
     "transform.test_branching_transform_op_terminator"()[^bb3] : () -> ()
   ^bb3:
     "transform.test_branching_transform_op_terminator"() : () -> ()
   }
   // expected-warning @below {{operand #0 may be used after free}}
-  transform.sequence %0 failures(propagate) {
+  transform.sequence %0 : !pdl.operation failures(propagate) {
     ^bb0(%arg0: !pdl.operation):
   }
   return
 }
 
 // -----
 
 func.func @use_after_free_recursive_side_effects() {
   transform.sequence failures(propagate) {
   ^bb0(%arg0: !pdl.operation):
     // expected-note @below {{allocated here}}
-    %0 = transform.sequence %arg0 failures(propagate) attributes { ord = 1 } {
+    %0 = transform.sequence %arg0 : !pdl.operation -> !pdl.operation failures(propagate) attributes { ord = 1 } {
     ^bb1(%arg1: !pdl.operation):
       yield %arg1 : !pdl.operation
-    } : !pdl.operation
-    transform.sequence %0 failures(propagate) attributes { ord = 2 } {
+    }
+    transform.sequence %0 : !pdl.operation failures(propagate) attributes { ord = 2 } {
     ^bb2(%arg2: !pdl.operation):
     }
-    transform.sequence %0 failures(propagate) attributes { ord = 3 } {
+    transform.sequence %0 : !pdl.operation failures(propagate) attributes { ord = 3 } {
     ^bb3(%arg3: !pdl.operation):
     }
     
     // `transform.sequence` has recursive side effects so it has the same "free"
     // as the child op it contains.
     // expected-note @below {{freed here}}
-    transform.sequence %0 failures(propagate) attributes { ord = 4 } {
+    transform.sequence %0 : !pdl.operation failures(propagate) attributes { ord = 4 } {
     ^bb4(%arg4: !pdl.operation):
       test_consume_operand_if_matches_param_or_fail %0[42]
     }
     // expected-warning @below {{operand #0 may be used after free}}
-    transform.sequence %0 failures(propagate) attributes { ord = 5 } {
+    transform.sequence %0 : !pdl.operation failures(propagate) attributes { ord = 5 } {
     ^bb3(%arg3: !pdl.operation):
     }
   }
   return
 }
 
 // -----
 
 func.func @use_after_free() {
   transform.sequence failures(propagate) {
   ^bb0(%arg0: !pdl.operation):
     // expected-note @below {{allocated here}}
-    %0 = transform.sequence %arg0 failures(propagate) attributes { ord = 1 } {
+    %0 = transform.sequence %arg0 : !pdl.operation -> !pdl.operation failures(propagate) attributes { ord = 1 } {
     ^bb1(%arg1: !pdl.operation):
       yield %arg1 : !pdl.operation
-    } : !pdl.operation
-    transform.sequence %0 failures(propagate) attributes { ord = 2 } {
+    }
+    transform.sequence %0 : !pdl.operation failures(propagate) attributes { ord = 2 } {
     ^bb2(%arg2: !pdl.operation):
     }
-    transform.sequence %0 failures(propagate) attributes { ord = 3 } {
+    transform.sequence %0 : !pdl.operation failures(propagate) attributes { ord = 3 } {
     ^bb3(%arg3: !pdl.operation):
     }
     
     // expected-note @below {{freed here}}
     test_consume_operand_if_matches_param_or_fail %0[42]
     // expected-warning @below {{operand #0 may be used after free}}
-    transform.sequence %0 failures(propagate) attributes { ord = 5 } {
+    transform.sequence %0 : !pdl.operation failures(propagate) attributes { ord = 5 } {
     ^bb3(%arg3: !pdl.operation):
     }
   }
   return
 }
 
 // -----
 
 // In the case of a control flow cycle, the operation that uses the value may
 // precede the one that frees it in the same block. Both operations should
 // be reported as use-after-free.
 func.func @use_after_free_self_cycle() {
   // expected-note @below {{allocated here}}
   %0 = transform.test_produce_param_or_forward_operand 42
   transform.test_transform_op_with_regions {
     "transform.test_branching_transform_op_terminator"() : () -> ()
   },
   {
   ^bb0:
     "transform.test_branching_transform_op_terminator"()[^bb1] : () -> ()
   ^bb1:
     // expected-warning @below {{operand #0 may be used after free}}
-    transform.sequence %0 failures(propagate) {
+    transform.sequence %0 : !pdl.operation failures(propagate) {
     ^bb0(%arg0: !pdl.operation):
     }
     // expected-warning @below {{operand #0 may be used after free}}
     // expected-note @below {{freed here}}
     transform.test_consume_operand_if_matches_param_or_fail %0[42]
     "transform.test_branching_transform_op_terminator"()[^bb1, ^bb2] : () -> ()
   ^bb2:
     "transform.test_branching_transform_op_terminator"() : () -> ()

mlir/test/Dialect/Transform/expensive-checks.mlir

 ^bb0(%arg0: !pdl.operation):
   pdl.pattern @return : benefit(1) {
     %0 = operands
     %1 = types
     %2 = operation "func.return"(%0 : !pdl.range<value>) -> (%1 : !pdl.range<type>)
     rewrite %2 with "transform.dialect"
   }
 
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     // expected-note @below {{handle to invalidated ops}}
-    %0 = pdl_match @return in %arg1
-    %1 = get_closest_isolated_parent %0
+    %0 = pdl_match @return in %arg1 : (!pdl.operation) -> !pdl.operation
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
     // expected-note @below {{invalidated by this transform op that consumes its operand #0}}
     test_consume_operand %1
     // expected-error @below {{op uses a handle invalidated by a previously executed transform op}}
     test_print_remark_at_operand %0, "remark"
   }
 }
 
 // -----
^bb0(%arg0: !pdl.operation):
   }
   pdl.pattern @return : benefit(1) {
     %0 = operands
     %1 = types
     %2 = operation "func.return"(%0 : !pdl.range<value>) -> (%1 : !pdl.range<type>)
     rewrite %2 with "transform.dialect"
   }
 
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
-    %0 = pdl_match @func in %arg1
-    %1 = pdl_match @return in %arg1
-    %2 = replicate num(%0) %1
+    %0 = pdl_match @func in %arg1 : (!pdl.operation) -> !pdl.operation
+    %1 = pdl_match @return in %arg1 : (!pdl.operation) -> !pdl.operation
+    %2 = replicate num(%0) %1 : !pdl.operation, !pdl.operation
     // expected-error @below {{a handle passed as operand #0 and consumed by this operation points to a payload operation more than once}}
     test_consume_operand %2
     test_print_remark_at_operand %0, "remark"
   }
 }
 
 
 // -----

mlir/test/Dialect/Transform/ops-invalid.mlir

 // RUN: mlir-opt %s -split-input-file -verify-diagnostics
 
-// expected-error @below {{expects the entry block to have one argument of type '!pdl.operation'}}
+// expected-error @below {{expects the entry block to have one argument of type implementing TransformTypeInterface}}
 transform.sequence failures(propagate) {
 }
 
 // -----
 
 // expected-note @below {{nested in another possible top-level op}}
 transform.sequence failures(propagate) {
 ^bb0(%arg0: !pdl.operation):

 transform.sequence failures(propagate) {
 ^bb0(%arg0: !pdl.operation):
   // expected-note @below {{op without interface}}
   arith.constant 42.0 : f32
 }
 
 // -----
 
-// expected-error @below {{expects the types of the terminator operands to match the types of the resul}}
-%0 = transform.sequence failures(propagate) {
+// expected-error @below {{expects the types of the terminator operands to match the types of the result}}
+%0 = transform.sequence -> !pdl.operation failures(propagate) {
 ^bb0(%arg0: !pdl.operation):
   // expected-note @below {{terminator}}
   transform.yield
-} : !pdl.operation
+}
 
 // -----
 
 // expected-note @below {{nested in another possible top-level op}}
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
   // expected-error @below {{expects the root operation to be provided for a nested op}}
   transform.sequence failures(propagate) {

 }
 
 // -----
 
 // expected-note @below {{parent operation}}
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
    // expected-error @below {{op cannot be nested}}
-  transform.with_pdl_patterns %arg0 {
+  transform.with_pdl_patterns %arg0 : !pdl.operation {
   ^bb1(%arg1: !pdl.operation):
   }
 }
 
 // -----
 
 // expected-error @below {{expects at least one region}}
 "transform.test_transform_unrestricted_op_no_interface"() : () -> ()

 
 transform.sequence failures(propagate) {
 ^bb0(%arg0: !pdl.operation):
   // expected-error @below {{result #0 has more than one potential consumer}}
   %0 = test_produce_param_or_forward_operand 42
   // expected-note @below {{used here as operand #0}}
   test_consume_operand_if_matches_param_or_fail %0[42]
   // expected-note @below {{used here as operand #0}}
-  transform.sequence %0 failures(propagate) {
+  transform.sequence %0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     test_consume_operand_if_matches_param_or_fail %arg1[42]
   }
 }
 
 // -----
 
 transform.sequence failures(propagate) {
 ^bb0(%arg0: !pdl.operation):
   // expected-error @below {{result #0 has more than one potential consumer}}
   %0 = test_produce_param_or_forward_operand 42
   // expected-note @below {{used here as operand #0}}
   test_consume_operand_if_matches_param_or_fail %0[42]
-  transform.sequence %0 failures(propagate) {
+  transform.sequence %0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     // expected-note @below {{used here as operand #0}}
     test_consume_operand_if_matches_param_or_fail %0[42]
   }
 }
 
 // -----
 
 transform.sequence failures(propagate) {
 ^bb0(%arg0: !pdl.operation):
   // expected-error @below {{result #0 has more than one potential consumer}}
   %0 = test_produce_param_or_forward_operand 42
   // expected-note @below {{used here as operand #0}}
   test_consume_operand_if_matches_param_or_fail %0[42]
   // expected-note @below {{used here as operand #0}}
-  transform.sequence %0 failures(propagate) {
+  transform.sequence %0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
-    transform.sequence %arg1 failures(propagate) {
+    transform.sequence %arg1 : !pdl.operation failures(propagate) {
     ^bb2(%arg2: !pdl.operation):
       test_consume_operand_if_matches_param_or_fail %arg2[42]
     }
   }
 }
 
 // -----
 
 transform.sequence failures(propagate) {
 ^bb1(%arg1: !pdl.operation):
   // expected-error @below {{expects at least one region}}
   transform.alternatives
 }
 
 // -----
 
 transform.sequence failures(propagate) {
 ^bb1(%arg1: !pdl.operation):
   // expected-error @below {{expects terminator operands to have the same type as results of the operation}}
-  %2 = transform.alternatives %arg1 -> !pdl.operation {
+  %2 = transform.alternatives %arg1 : !pdl.operation -> !pdl.operation {
   ^bb2(%arg2: !pdl.operation):
     transform.yield %arg2 : !pdl.operation
   }, {
   ^bb2(%arg2: !pdl.operation):
     // expected-note @below {{terminator}}
     transform.yield
   }
 }
 
 // -----
 
-// expected-error @below {{expects the entry block to have one argument of type '!pdl.operation'}}
+// expected-error @below {{expects the entry block to have one argument of type implementing TransformTypeInterface}}
 transform.alternatives {
 ^bb0:
   transform.yield
 }
 
 // -----
 
 transform.sequence failures(propagate) {
 ^bb0(%arg0: !pdl.operation):
   // expected-error @below {{result #0 has more than one potential consumer}}
   %0 = test_produce_param_or_forward_operand 42
   // expected-note @below {{used here as operand #0}}
-  transform.foreach %0 {
+  transform.foreach %0 : !pdl.operation {
   ^bb1(%arg1: !pdl.operation):
     transform.test_consume_operand %arg1
   }
   // expected-note @below {{used here as operand #0}}
   transform.test_consume_operand %0
 }

mlir/test/Dialect/Transform/ops.mlir

 // RUN: mlir-opt %s | mlir-opt | FileCheck %s
 
 // CHECK: transform.sequence
 // CHECK: ^{{.+}}(%{{.+}}: !pdl.operation):
 transform.sequence failures(propagate) {
 ^bb0(%arg0: !pdl.operation):
-  // CHECK: sequence %{{.+}}
+  // CHECK: sequence %{{.+}} : !pdl.operation
   // CHECK: ^{{.+}}(%{{.+}}: !pdl.operation):
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
   }
 }
 
 // CHECK: transform.with_pdl_patterns
 // CHECK: ^{{.+}}(%[[ARG:.+]]: !pdl.operation):
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  // CHECK: sequence %[[ARG]]
-  sequence %arg0 failures(propagate) {
+  // CHECK: sequence %[[ARG]] : !pdl.operation
+  sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
   }
 }
 
 // CHECK: transform.sequence
 // CHECK: ^{{.+}}(%[[ARG:.+]]: !pdl.operation):
 transform.sequence failures(propagate) {
 ^bb0(%arg0: !pdl.operation):
-  // CHECK: with_pdl_patterns %[[ARG]]
-  with_pdl_patterns %arg0 {
+  // CHECK: with_pdl_patterns %[[ARG]] : !pdl.operation
+  with_pdl_patterns %arg0 : !pdl.operation {
   ^bb1(%arg1: !pdl.operation):
   }
 }
 
 // Using the same value multiple times without consuming it is fine.
 // CHECK: transform.sequence
-// CHECK: %[[V:.+]] = sequence
+// CHECK: %[[V:.+]] = sequence %{{.*}} : !pdl.operation -> !pdl.operation
 // CHECK: sequence %[[V]]
 // CHECK: sequence %[[V]]
 transform.sequence failures(propagate) {
 ^bb0(%arg0: !pdl.operation):
-  %0 = transform.sequence %arg0 failures(propagate) {
+  %0 = transform.sequence %arg0 : !pdl.operation -> !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     yield %arg1 : !pdl.operation
-  } : !pdl.operation
-  transform.sequence %0 failures(propagate) {
+  }
+  transform.sequence %0 : !pdl.operation failures(propagate) {
   ^bb2(%arg2: !pdl.operation):
   }
-  transform.sequence %0 failures(propagate) {
+  transform.sequence %0 : !pdl.operation failures(propagate) {
   ^bb3(%arg3: !pdl.operation):
   }
 }
 
 // CHECK: transform.sequence
 // CHECK: foreach
 transform.sequence failures(propagate) {
 ^bb0(%arg0: !pdl.operation):
-  transform.foreach %arg0 {
+  transform.foreach %arg0 : !pdl.operation {
   ^bb1(%arg1: !pdl.operation):
   }
 }
 
 // CHECK: transform.sequence
 transform.sequence failures(propagate) {
 ^bb0(%arg0: !pdl.operation):
   // CHECK: cast %{{.*}} : !pdl.operation to !transform.any_op
   %0 = cast %arg0: !pdl.operation to !transform.any_op
 }

mlir/test/Dialect/Transform/selective-targeting.mlir

pdl.pattern @pdl_target_attrC : benefit(1) {
     %args = operands
     %results = types
     %attr = attribute
     %0 = operation "linalg.matmul"(%args : !pdl.range<value>) {"test.attrC" = %attr}-> (%results : !pdl.range<type>)
     // TODO: we don't want this, but it is the required terminator for pdl.pattern
     rewrite %0 with "transform.dialect"
   }
 
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
-    %0 = pdl_match @pdl_target_attrA in %arg1
+    %0 = pdl_match @pdl_target_attrA in %arg1 : (!pdl.operation) -> !pdl.operation
     transform.structured.tile %0 [4, 4, 4]
-    %1 = pdl_match @pdl_target_attrC in %arg1
-    %2 = transform.get_closest_isolated_parent %1
+    %1 = pdl_match @pdl_target_attrC in %arg1 : (!pdl.operation) -> !pdl.operation
+    %2 = transform.get_closest_isolated_parent %1 : (!pdl.operation) -> !pdl.operation
     transform.structured.vectorize %2
   }
 }
 
 // -----
 
 // CHECK-LABEL: @vectorize_one
 func.func @vectorize_one(
pdl.pattern @pdl_target : benefit(1) {
     %args = operands
     %results = types
     %attr = attribute
     %0 = operation "linalg.matmul"(%args : !pdl.range<value>) {"test.attrA" = %attr}-> (%results : !pdl.range<type>)
     // TODO: we don't want this, but it is the required terminator for pdl.pattern
     rewrite %0 with "transform.dialect"
   }
 
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
-    %0 = pdl_match @pdl_target in %arg1
-    %1 = get_closest_isolated_parent %0
+    %0 = pdl_match @pdl_target in %arg1 : (!pdl.operation) -> !pdl.operation
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
     transform.structured.vectorize %1
   }
 }
 
 // -----
 
 // CHECK-LABEL: @vectorize_all
 func.func @vectorize_all(

mlir/test/Dialect/Transform/test-interpreter.mlir

 %1 = transform.test_copy_payload %0
 // expected-remark @below {{succeeded}}
 transform.test_consume_operand_if_matches_param_or_fail %0[42]
 
 // -----
 
 transform.sequence failures(propagate) {
 ^bb0(%arg0: !pdl.operation):
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb0(%arg1: !pdl.operation):
     // expected-remark @below {{applying transformation "a"}}
     test_transform_op "a"
     // expected-remark @below {{applying transformation "b"}}
     test_transform_op "b"
     // expected-remark @below {{applying transformation "c"}}
     test_transform_op "c"
   }
   // expected-remark @below {{applying transformation "d"}}
   test_transform_op "d"
   // expected-remark @below {{applying transformation "e"}}
   test_transform_op "e"
 }
 
 // -----
 
 transform.sequence failures(propagate) {
 ^bb0(%arg0: !pdl.operation):
   %0 = test_produce_param_or_forward_operand 42
-  sequence %0 failures(propagate) {
+  sequence %0 : !pdl.operation failures(propagate) {
   ^bb0(%arg1: !pdl.operation):
     // expected-remark @below {{succeeded}}
     test_consume_operand_if_matches_param_or_fail %arg1[42]
   }
 }
 
 // -----
 
 transform.sequence failures(propagate) {
 ^bb0(%arg0: !pdl.operation):
-  %0 = sequence %arg0 failures(propagate) {
+  %0 = sequence %arg0 : !pdl.operation -> !pdl.operation failures(propagate) {
   ^bb0(%arg1: !pdl.operation):
     %1 = test_produce_param_or_forward_operand 42
     yield %1 : !pdl.operation
-  } : !pdl.operation
+  }
   // expected-remark @below {{succeeded}}
   test_consume_operand_if_matches_param_or_fail %0[42]
 }
 
 // -----
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb0(%arg1: !pdl.operation):
-    %0 = pdl_match @some in %arg1
+    %0 = pdl_match @some in %arg1 : (!pdl.operation) -> !pdl.operation
     test_print_remark_at_operand %0, "matched"
   }
 
   pdl.pattern @some : benefit(1) {
     %0 = pdl.operation "test.some_op"
     pdl.rewrite %0 with "transform.dialect"
   }
 

 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
   pdl.pattern @const : benefit(1) {
     %r = pdl.types
     %0 = pdl.operation "arith.constant" -> (%r : !pdl.range<type>)
     pdl.rewrite %0 with "transform.dialect"
   }
 
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
-    %f = pdl_match @const in %arg1
+    %f = pdl_match @const in %arg1 : (!pdl.operation) -> !pdl.operation
     // CHECK: %{{.+}} = get_closest_isolated_parent %{{.+}}
-    %m = get_closest_isolated_parent %f
+    %m = get_closest_isolated_parent %f : (!pdl.operation) -> !pdl.operation
     test_print_remark_at_operand %m, "parent function"
   }
 }
 
 // -----
 
 func.func @foo() {
   %0 = arith.constant 0 : i32
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
   pdl.pattern @match_func : benefit(1) {
     %0 = pdl.operands
     %1 = pdl.types
     %2 = pdl.operation "func.func"(%0 : !pdl.range<value>) -> (%1 : !pdl.range<type>)
     pdl.rewrite %2 with "transform.dialect"
   }
 
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     // This is necessary to run the transformation on something other than the
     // top-level module, "alternatives" cannot be run on that.
-    %0 = pdl_match @match_func in %arg1
-    transform.alternatives %0 {
+    %0 = pdl_match @match_func in %arg1 : (!pdl.operation) -> !pdl.operation
+    transform.alternatives %0 : !pdl.operation {
     ^bb2(%arg2: !pdl.operation):
       %1 = transform.test_produce_param_or_forward_operand 42
       // This operation fails, which triggers the next alternative without
       // reporting the error.
       transform.test_consume_operand_if_matches_param_or_fail %1[43]
     }, {
     ^bb2(%arg2: !pdl.operation):
       %1 = transform.test_produce_param_or_forward_operand 42

 ^bb0(%arg0: !pdl.operation):
   pdl.pattern @match_call : benefit(1) {
     %0 = pdl.operands
     %1 = pdl.types
     %2 = pdl.operation "func.call"(%0 : !pdl.range<value>) -> (%1 : !pdl.range<type>)
     pdl.rewrite %2 with "transform.dialect"
   }
 
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
-    %0 = pdl_match @match_call in %arg1
-    %1 = get_closest_isolated_parent %0
+    %0 = pdl_match @match_call in %arg1 : (!pdl.operation) -> !pdl.operation
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
     // expected-error @below {{all alternatives failed}}
-    transform.alternatives %1 {
+    transform.alternatives %1 : !pdl.operation {
     ^bb2(%arg2: !pdl.operation):
-      %2 = transform.pdl_match @match_call in %arg2
+      %2 = transform.pdl_match @match_call in %arg2 : (!pdl.operation) -> !pdl.operation
       // expected-remark @below {{applying}}
       transform.test_emit_remark_and_erase_operand %2, "applying" {fail_after_erase}
     }
   }
 }
 
 // -----
 

 ^bb0(%arg0: !pdl.operation):
   pdl.pattern @match_call : benefit(1) {
     %0 = pdl.operands
     %1 = pdl.types
     %2 = pdl.operation "func.call"(%0 : !pdl.range<value>) -> (%1 : !pdl.range<type>)
     pdl.rewrite %2 with "transform.dialect"
   }
 
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
-    %0 = pdl_match @match_call in %arg1
-    %1 = get_closest_isolated_parent %0
-    transform.alternatives %1 {
+    %0 = pdl_match @match_call in %arg1 : (!pdl.operation) -> !pdl.operation
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
+    transform.alternatives %1 : !pdl.operation {
     ^bb2(%arg2: !pdl.operation):
-      %2 = transform.pdl_match @match_call in %arg2
+      %2 = transform.pdl_match @match_call in %arg2 : (!pdl.operation) -> !pdl.operation
       // expected-remark @below {{applying}}
       transform.test_emit_remark_and_erase_operand %2, "applying" {fail_after_erase}
     }, {
     ^bb2(%arg2: !pdl.operation):
-      %2 = transform.pdl_match @match_call in %arg2
+      %2 = transform.pdl_match @match_call in %arg2 : (!pdl.operation) -> !pdl.operation
       transform.test_print_remark_at_operand %2, "still here"
       // This alternative succeeds.
     }, {
     ^bb2(%arg2: !pdl.operation):
       // This alternative is never run, so we must not have a remark here.
-      %2 = transform.pdl_match @match_call in %arg2
+      %2 = transform.pdl_match @match_call in %arg2 : (!pdl.operation) -> !pdl.operation
       transform.test_emit_remark_and_erase_operand %2, "should not happen" {fail_after_erase}
     }
   }
 }
 
 // -----
 
 func.func private @bar()

 ^bb0(%arg0: !pdl.operation):
   pdl.pattern @match_call : benefit(1) {
     %0 = pdl.operands
     %1 = pdl.types
     %2 = pdl.operation "func.call"(%0 : !pdl.range<value>) -> (%1 : !pdl.range<type>)
     pdl.rewrite %2 with "transform.dialect"
   }
 
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
-    %0 = pdl_match @match_call in %arg1
-    %1 = get_closest_isolated_parent %0
-    transform.alternatives %1 {
+    %0 = pdl_match @match_call in %arg1 : (!pdl.operation) -> !pdl.operation
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
+    transform.alternatives %1 : !pdl.operation {
     ^bb2(%arg2: !pdl.operation):
-      %2 = transform.pdl_match @match_call in %arg2
+      %2 = transform.pdl_match @match_call in %arg2 : (!pdl.operation) -> !pdl.operation
       // expected-remark @below {{applying}}
       transform.test_emit_remark_and_erase_operand %2, "applying" {fail_after_erase}
     }, {
     ^bb2(%arg2: !pdl.operation):
-      %2 = transform.pdl_match @match_call in %arg2
+      %2 = transform.pdl_match @match_call in %arg2 : (!pdl.operation) -> !pdl.operation
       // expected-remark @below {{applying second time}}
       transform.test_emit_remark_and_erase_operand %2, "applying second time"
     }
   }
 }
 
 // -----
 
 func.func private @bar()
 
 func.func @foo() {
   call @bar() : () -> ()
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
   pdl.pattern @match_call : benefit(1) {
     %0 = pdl.operands
     %1 = pdl.types
     %2 = pdl.operation "func.call"(%0 : !pdl.range<value>) -> (%1 : !pdl.range<type>)
     pdl.rewrite %2 with "transform.dialect"
   }
 
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
-    %0 = pdl_match @match_call in %arg1
-    %1 = get_closest_isolated_parent %0
-    %2 = transform.alternatives %1 -> !pdl.operation {
+    %0 = pdl_match @match_call in %arg1 : (!pdl.operation) -> !pdl.operation
+    %1 = get_closest_isolated_parent %0 : (!pdl.operation) -> !pdl.operation
+    %2 = transform.alternatives %1 : !pdl.operation -> !pdl.operation {
     ^bb2(%arg2: !pdl.operation):
-      %3 = transform.pdl_match @match_call in %arg2
+      %3 = transform.pdl_match @match_call in %arg2 : (!pdl.operation) -> !pdl.operation
       // expected-remark @below {{applying}}
       transform.test_emit_remark_and_erase_operand %3, "applying" {fail_after_erase}
       %4 = transform.test_produce_param_or_forward_operand 43
       transform.yield %4 : !pdl.operation
     }, {
     ^bb2(%arg2: !pdl.operation):
       %4 = transform.test_produce_param_or_forward_operand 42
       transform.yield %4 : !pdl.operation
func.func @bar() {
     %0 = arith.constant 0 : i32
     %1 = arith.constant 1 : i32
     return
   }
 
   transform.sequence failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     // expected-error @below {{scope must not contain the transforms being applied}}
-    transform.alternatives %arg1 {
+    transform.alternatives %arg1 : !pdl.operation {
     ^bb2(%arg2: !pdl.operation):
       %0 = transform.test_produce_param_or_forward_operand 42
       transform.test_consume_operand_if_matches_param_or_fail %0[43]
     }, {
     ^bb2(%arg2: !pdl.operation):
       %0 = transform.test_produce_param_or_forward_operand 42
       transform.test_consume_operand_if_matches_param_or_fail %0[42]
     }
^bb0(%arg0: !pdl.operation):
   pdl.pattern @match_const : benefit(1) {
     %0 = pdl.operands
     %1 = pdl.types
     %2 = pdl.operation "arith.constant"(%0 : !pdl.range<value>) -> (%1 : !pdl.range<type>)
     pdl.rewrite %2 with "transform.dialect"
   }
 
 
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
-    %0 = transform.pdl_match @match_const in %arg1
+    %0 = transform.pdl_match @match_const in %arg1 : (!pdl.operation) -> !pdl.operation
     %1 = transform.loop.get_parent_for %0
     // expected-error @below {{only isolated-from-above ops can be alternative scopes}}
-    alternatives %1 {
+    alternatives %1 : !pdl.operation {
     ^bb2(%arg2: !pdl.operation):
     }
   }
 }
 // -----
 
 func.func @foo() {
   // expected-note @below {{when applied to this op}}
   "op" () : () -> ()
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
   pdl.pattern @some : benefit(1) {
     %0 = pdl.operands
     %1 = pdl.types
     %2 = pdl.operation "op"(%0 : !pdl.range<value>) -> (%1 : !pdl.range<type>)
     pdl.rewrite %2 with "transform.dialect"
   }
 
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb0(%arg1: !pdl.operation):
-    %0 = pdl_match @some in %arg1
+    %0 = pdl_match @some in %arg1 : (!pdl.operation) -> !pdl.operation
     // expected-error @below {{applications of transform.test_wrong_number_of_results expected to produce 3 results (actually produced 1).}}
     // expected-note @below {{If you need variadic results, consider a generic `apply` instead of the specialized `applyToOne`.}}
     // expected-note @below {{Producing 3 null results is allowed if the use case warrants it.}}
     transform.test_wrong_number_of_results %0
   }
 }
 
 // -----

 ^bb0(%arg0: !pdl.operation):
   pdl.pattern @some : benefit(1) {
     %0 = pdl.operands
     %1 = pdl.types
     %2 = pdl.operation "op"(%0 : !pdl.range<value>) -> (%1 : !pdl.range<type>)
     pdl.rewrite %2 with "transform.dialect"
   }
 
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb0(%arg1: !pdl.operation):
-    %0 = pdl_match @some in %arg1
+    %0 = pdl_match @some in %arg1 : (!pdl.operation) -> !pdl.operation
     // expected-error @below {{applications of transform.test_wrong_number_of_multi_results expected to produce 1 results (actually produced 0)}}
     // expected-note @below {{If you need variadic results, consider a generic `apply` instead of the specialized `applyToOne`.}}
     // expected-note @below {{Producing 1 null results is allowed if the use case warrants it.}}
     transform.test_wrong_number_of_multi_results %0
   }
 }
 
 // -----

 ^bb0(%arg0: !pdl.operation):
   pdl.pattern @some : benefit(1) {
     %0 = pdl.operands
     %1 = pdl.types
     %2 = pdl.operation "op"(%0 : !pdl.range<value>) -> (%1 : !pdl.range<type>)
     pdl.rewrite %2 with "transform.dialect"
   }
 
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb0(%arg1: !pdl.operation):
-    %0 = pdl_match @some in %arg1
+    %0 = pdl_match @some in %arg1 : (!pdl.operation) -> !pdl.operation
     // Transform matches 3 ops and produces 2 results.
     %1:2 = transform.test_correct_number_of_multi_results %0
   }
 }
 
 // -----
 
 func.func @foo() {
   "wrong_op_name" () : () -> ()
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
   pdl.pattern @some : benefit(1) {
     %0 = pdl.operands
     %1 = pdl.types
     %2 = pdl.operation "op"(%0 : !pdl.range<value>) -> (%1 : !pdl.range<type>)
     pdl.rewrite %2 with "transform.dialect"
   }
 
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb0(%arg1: !pdl.operation):
-    %0 = pdl_match @some in %arg1
+    %0 = pdl_match @some in %arg1 : (!pdl.operation) -> !pdl.operation
     // Transform fails to match any but still produces 2 results.
     %1:2 = transform.test_correct_number_of_multi_results %0
   }
 }
 
 // -----
 
 func.func @foo() {
   // expected-note @below {{when applied to this op}}
   "op" () : () -> ()
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
   pdl.pattern @some : benefit(1) {
     %0 = pdl.operands
     %1 = pdl.types
     %2 = pdl.operation "op"(%0 : !pdl.range<value>) -> (%1 : !pdl.range<type>)
     pdl.rewrite %2 with "transform.dialect"
   }
 
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb0(%arg1: !pdl.operation):
-    %0 = pdl_match @some in %arg1
+    %0 = pdl_match @some in %arg1 : (!pdl.operation) -> !pdl.operation
     // expected-error @below {{unexpected application of transform.test_mixed_null_and_non_null_results produces both null and non null results.}}
     transform.test_mixed_null_and_non_null_results %0
   }
 }
 
 // -----
 
 // Expecting to match all operations by merging the handles that matched addi
^bb0(%arg0: !pdl.operation):
   }
   pdl.pattern @subi : benefit(1) {
     %0 = pdl.operands
     %1 = pdl.types
     %2 = pdl.operation "arith.subi"(%0 : !pdl.range<value>) -> (%1 : !pdl.range<type>)
     pdl.rewrite %2 with "transform.dialect"
   }
 
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb0(%arg1: !pdl.operation):
-    %0 = pdl_match @addi in %arg1
-    %1 = pdl_match @subi in %arg1
-    %2 = merge_handles %0, %1
+    %0 = pdl_match @addi in %arg1 : (!pdl.operation) -> !pdl.operation
+    %1 = pdl_match @subi in %arg1 : (!pdl.operation) -> !pdl.operation
+    %2 = merge_handles %0, %1 : !pdl.operation
     test_print_remark_at_operand %2, "matched"
   }
 }
 
 // -----
 
 func.func @foo() {
   "op" () { target_me } : () -> ()
   // expected-note @below {{when applied to this op}}
   "op" () : () -> ()
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
   pdl.pattern @some : benefit(1) {
     %0 = pdl.operands
     %1 = pdl.types
     %2 = pdl.operation "op"(%0 : !pdl.range<value>) -> (%1 : !pdl.range<type>)
     pdl.rewrite %2 with "transform.dialect"
   }
 
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb0(%arg1: !pdl.operation):
-    %0 = pdl_match @some in %arg1
+    %0 = pdl_match @some in %arg1 : (!pdl.operation) -> !pdl.operation
     // expected-error @below {{failed to apply}}
     transform.test_mixed_sucess_and_silenceable %0
   }
 }
 
 // -----
 
 func.func @foo() {
   "op" () : () -> ()
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
   pdl.pattern @some : benefit(1) {
     %0 = pdl.operands
     %1 = pdl.types
     %2 = pdl.operation "op"(%0 : !pdl.range<value>) -> (%1 : !pdl.range<type>)
     pdl.rewrite %2 with "transform.dialect"
   }
 
-  transform.sequence %arg0 failures(suppress) {
+  transform.sequence %arg0 : !pdl.operation failures(suppress) {
   ^bb0(%arg1: !pdl.operation):
-    %0 = pdl_match @some in %arg1
+    %0 = pdl_match @some in %arg1 : (!pdl.operation) -> !pdl.operation
     // Not expecting error here because we are suppressing it.
     // expected-remark @below {{foo}}
     test_emit_remark_and_erase_operand %0, "foo" {fail_after_erase}
   }
 }
 
 // -----
 
 func.func @foo() {
   "op" () : () -> ()
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
   pdl.pattern @some : benefit(1) {
     %0 = pdl.operands
     %1 = pdl.types
     %2 = pdl.operation "op"(%0 : !pdl.range<value>) -> (%1 : !pdl.range<type>)
     pdl.rewrite %2 with "transform.dialect"
   }
 
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb0(%arg1: !pdl.operation):
-    %0 = pdl_match @some in %arg1
+    %0 = pdl_match @some in %arg1 : (!pdl.operation) -> !pdl.operation
     // expected-error @below {{silenceable error}}
     // expected-remark @below {{foo}}
     test_emit_remark_and_erase_operand %0, "foo" {fail_after_erase}
   }
 }
 
 
 // -----
 
 module {
   func.func private @foo()
   func.func private @bar()
 
   transform.with_pdl_patterns {
   ^bb0(%arg0: !pdl.operation):
     pdl.pattern @func : benefit(1) {
       %0 = pdl.operands
       %1 = pdl.types
       %2 = pdl.operation "func.func"(%0 : !pdl.range<value>) -> (%1 : !pdl.range<type>)
       pdl.rewrite %2 with "transform.dialect"
     }
 
-    transform.sequence %arg0 failures(propagate) {
+    transform.sequence %arg0 : !pdl.operation failures(propagate) {
     ^bb0(%arg1: !pdl.operation):
-      %0 = pdl_match @func in %arg1
-      %1 = replicate num(%0) %arg1
+      %0 = pdl_match @func in %arg1 : (!pdl.operation) -> !pdl.operation
+      %1 = replicate num(%0) %arg1 : !pdl.operation, !pdl.operation
       // expected-remark @below {{2}}
       test_print_number_of_associated_payload_ir_ops %1
-      %2 = replicate num(%0) %1
+      %2 = replicate num(%0) %1 : !pdl.operation, !pdl.operation
       // expected-remark @below {{4}}
       test_print_number_of_associated_payload_ir_ops %2
     }
   }
 }
 
 // -----
 
 func.func @bar() {
   // expected-remark @below {{transform applied}}
   %0 = arith.constant 0 : i32
   // expected-remark @below {{transform applied}}
   %1 = arith.constant 1 : i32
   return
 }
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
   pdl.pattern @const : benefit(1) {
     %r = pdl.types
     %0 = pdl.operation "arith.constant" -> (%r : !pdl.range<type>)
     pdl.rewrite %0 with "transform.dialect"
   }
 
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
-    %f = pdl_match @const in %arg1
-    transform.foreach %f {
+    %f = pdl_match @const in %arg1 : (!pdl.operation) -> !pdl.operation
+    transform.foreach %f : !pdl.operation {
     ^bb2(%arg2: !pdl.operation):
       // expected-remark @below {{1}}
       transform.test_print_number_of_associated_payload_ir_ops %arg2
       transform.test_print_remark_at_operand %arg2, "transform applied"
     }
   }
 }
 
^bb0(%arg0: !pdl.operation):
   }
 
   pdl.pattern @execute_region : benefit(1) {
     %r = pdl.types
     %0 = pdl.operation "scf.execute_region" -> (%r : !pdl.range<type>)
     pdl.rewrite %0 with "transform.dialect"
   }
 
-  transform.sequence %arg0 failures(propagate) {
+  transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
-    %f = pdl_match @execute_region in %arg1
-    %results = transform.foreach %f -> !pdl.operation {
+    %f = pdl_match @execute_region in %arg1 : (!pdl.operation) -> !pdl.operation
+    %results = transform.foreach %f : !pdl.operation -> !pdl.operation {
     ^bb2(%arg2: !pdl.operation):
-      %g = transform.pdl_match @const in %arg2
+      %g = transform.pdl_match @const in %arg2 : (!pdl.operation) -> !pdl.operation
       transform.yield %g : !pdl.operation
     }
 
     // expected-remark @below {{3}}
     transform.test_print_number_of_associated_payload_ir_ops %results
     transform.test_print_remark_at_operand %results, "transform applied"
   }
 }
 
 // -----
 
 func.func @get_parent_for_op_no_loop(%arg0: index, %arg1: index) {
   // expected-remark @below {{found muli}}
   %0 = arith.muli %arg0, %arg1 : index  
   arith.addi %0, %arg1 : index  
   return
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%arg1: !pdl.operation):
   %addi = transform.structured.match ops{["arith.addi"]} in %arg1
-  %muli = get_producer_of_operand %addi[0]
+  %muli = get_producer_of_operand %addi[0] : (!pdl.operation) -> !pdl.operation
   transform.test_print_remark_at_operand %muli, "found muli"
 }
 
 // -----
 
 func.func @get_parent_for_op_no_loop(%arg0: index, %arg1: index) {
   // expected-note @below {{target op}}
   %0 = arith.muli %arg0, %arg1 : index  
   return
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%arg1: !pdl.operation):
   %muli = transform.structured.match ops{["arith.muli"]} in %arg1
   // expected-error @below {{could not find a producer for operand number: 0 of}}
-  %bbarg = get_producer_of_operand %muli[0]
+  %bbarg = get_producer_of_operand %muli[0] : (!pdl.operation) -> !pdl.operation
 
 }
 
 // -----
 
 func.func @split_handles(%a: index, %b: index, %c: index) {
   %0 = arith.muli %a, %b : index  
   %1 = arith.muli %a, %c : index  
   return
 }
 
 transform.sequence failures(propagate) {
 ^bb1(%fun: !pdl.operation):
   %muli = transform.structured.match ops{["arith.muli"]} in %fun
-  %h:2 = split_handles %muli in [2]
+  %h:2 = split_handles %muli in [2] : (!pdl.operation) -> (!pdl.operation, !pdl.operation)
   // expected-remark @below {{1}}
   transform.test_print_number_of_associated_payload_ir_ops %h#0
   %muli_2 = transform.structured.match ops{["arith.muli"]} in %fun
   // expected-error @below {{expected to contain 3 operation handles but it only contains 2 handles}}
-  %h_2:3 = split_handles %muli_2 in [3]
+  %h_2:3 = split_handles %muli_2 in [3] : (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation)
 }
 
 // -----
 
 func.func @split_handles(%a: index, %b: index, %c: index) {
   %0 = arith.muli %a, %b : index  
   %1 = arith.muli %a, %c : index  
   return
 }
 
 transform.sequence failures(suppress) {
 ^bb1(%fun: !pdl.operation):
   %muli = transform.structured.match ops{["arith.muli"]} in %fun
-  %h:2 = split_handles %muli in [2]
+  %h:2 = split_handles %muli in [2] : (!pdl.operation) -> (!pdl.operation, !pdl.operation)
   // expected-remark @below {{1}}
   transform.test_print_number_of_associated_payload_ir_ops %h#0
   %muli_2 = transform.structured.match ops{["arith.muli"]} in %fun
   // Silenceable failure and all handles are now empty.
-  %h_2:3 = split_handles %muli_2 in [3]
+  %h_2:3 = split_handles %muli_2 in [3] : (!pdl.operation) -> (!pdl.operation, !pdl.operation, !pdl.operation)
   // expected-remark @below {{0}}
   transform.test_print_number_of_associated_payload_ir_ops %h_2#0
 }
 
 // -----
 
 "test.some_op"() : () -> ()
 "other_dialect.other_op"() : () -> ()
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
   pdl.pattern @some : benefit(1) {
     %0 = pdl.operation "test.some_op"
     pdl.rewrite %0 with "transform.dialect"
   }
 
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
-    %0 = pdl_match @some in %arg1
+    %0 = pdl_match @some in %arg1 : (!pdl.operation) -> !pdl.operation
     %2 = transform.cast %0 : !pdl.operation to !transform.test_dialect_op
     transform.cast %2 : !transform.test_dialect_op to !pdl.operation
   }
 }
 
 // -----
 
 "test.some_op"() : () -> ()
 "other_dialect.other_op"() : () -> ()
 
 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
   pdl.pattern @other : benefit(1) {
     %0 = pdl.operation "other_dialect.other_op"
     pdl.rewrite %0 with "transform.dialect"
   }
 
-  sequence %arg0 failures(propagate) {
+  sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
-    %0 = pdl_match @other in %arg1
+    %0 = pdl_match @other in %arg1 : (!pdl.operation) -> !pdl.operation
     // expected-error @below {{expected the payload operation to belong to the 'test' dialect}}
     %2 = transform.cast %0 : !pdl.operation to !transform.test_dialect_op
     transform.cast %2 : !transform.test_dialect_op to !pdl.operation
   }
 }

mlir/test/python/dialects/transform.py

with InsertionPoint(module.body):
       f()
     print(module)
   return f
 
 
 @run
 def testSequenceOp():
   sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE,
-                                  [pdl.OperationType.get()])
+                                  [pdl.OperationType.get()],
+                                  pdl.OperationType.get())
   with InsertionPoint(sequence.body):
     transform.YieldOp([sequence.bodyTarget])
   # CHECK-LABEL: TEST: testSequenceOp
-  # CHECK: = transform.sequence failures(propagate) {
+  # CHECK: = transform.sequence -> !pdl.operation failures(propagate) {
   # CHECK: ^{{.*}}(%[[ARG0:.+]]: !pdl.operation):
   # CHECK:   yield %[[ARG0]] : !pdl.operation
-  # CHECK: } : !pdl.operation
+  # CHECK: }
 
 
 @run
 def testNestedSequenceOp():
-  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE)
+  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE, [], pdl.OperationType.get())
   with InsertionPoint(sequence.body):
-    nested = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE, sequence.bodyTarget)
+    nested = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE, [], sequence.bodyTarget)
     with InsertionPoint(nested.body):
       doubly_nested = transform.SequenceOp(
           transform.FailurePropagationMode.PROPAGATE,
           [pdl.OperationType.get()], nested.bodyTarget)
       with InsertionPoint(doubly_nested.body):
         transform.YieldOp([doubly_nested.bodyTarget])
       transform.YieldOp()
     transform.YieldOp()
   # CHECK-LABEL: TEST: testNestedSequenceOp
   # CHECK: transform.sequence failures(propagate) {
   # CHECK: ^{{.*}}(%[[ARG0:.+]]: !pdl.operation):
-  # CHECK:   sequence %[[ARG0]] failures(propagate) {
+  # CHECK:   sequence %[[ARG0]] : !pdl.operation failures(propagate) {
   # CHECK:   ^{{.*}}(%[[ARG1:.+]]: !pdl.operation):
-  # CHECK:     = sequence %[[ARG1]] failures(propagate) {
+  # CHECK:     = sequence %[[ARG1]] : !pdl.operation -> !pdl.operation failures(propagate) {
   # CHECK:     ^{{.*}}(%[[ARG2:.+]]: !pdl.operation):
   # CHECK:       yield %[[ARG2]] : !pdl.operation
-  # CHECK:     } : !pdl.operation
+  # CHECK:     }
   # CHECK:   }
   # CHECK: }
 
 
 @run
 def testTransformPDLOps():
-  withPdl = transform.WithPDLPatternsOp()
+  withPdl = transform.WithPDLPatternsOp(pdl.OperationType.get())
   with InsertionPoint(withPdl.body):
     sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE,
                                     [pdl.OperationType.get()],
                                     withPdl.bodyTarget)
     with InsertionPoint(sequence.body):
-      match = transform.PDLMatchOp(sequence.bodyTarget, "pdl_matcher")
+      match = transform.PDLMatchOp(pdl.OperationType.get(), sequence.bodyTarget, "pdl_matcher")
       transform.YieldOp(match)
   # CHECK-LABEL: TEST: testTransformPDLOps
   # CHECK: transform.with_pdl_patterns {
   # CHECK: ^{{.*}}(%[[ARG0:.+]]: !pdl.operation):
-  # CHECK:   = sequence %[[ARG0]] failures(propagate) {
+  # CHECK:   = sequence %[[ARG0]] : !pdl.operation -> !pdl.operation failures(propagate) {
   # CHECK:   ^{{.*}}(%[[ARG1:.+]]: !pdl.operation):
   # CHECK:     %[[RES:.+]] = pdl_match @pdl_matcher in %[[ARG1]]
   # CHECK:     yield %[[RES]] : !pdl.operation
-  # CHECK:   } : !pdl.operation
+  # CHECK:   }
   # CHECK: }
 
 
 @run
 def testGetClosestIsolatedParentOp():
-  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE)
+  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE, [], pdl.OperationType.get())
   with InsertionPoint(sequence.body):
-    transform.GetClosestIsolatedParentOp(sequence.bodyTarget)
+    transform.GetClosestIsolatedParentOp(pdl.OperationType.get(), sequence.bodyTarget)
     transform.YieldOp()
   # CHECK-LABEL: TEST: testGetClosestIsolatedParentOp
   # CHECK: transform.sequence
   # CHECK: ^{{.*}}(%[[ARG1:.+]]: !pdl.operation):
   # CHECK:   = get_closest_isolated_parent %[[ARG1]]
 
 
 @run
 def testMergeHandlesOp():
-  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE)
+  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE, [], pdl.OperationType.get())
   with InsertionPoint(sequence.body):
     transform.MergeHandlesOp([sequence.bodyTarget])
     transform.YieldOp()
   # CHECK-LABEL: TEST: testMergeHandlesOp
   # CHECK: transform.sequence
   # CHECK: ^{{.*}}(%[[ARG1:.+]]: !pdl.operation):
   # CHECK:   = merge_handles %[[ARG1]]
 
 
 @run
 def testReplicateOp():
-  with_pdl = transform.WithPDLPatternsOp()
+  with_pdl = transform.WithPDLPatternsOp(pdl.OperationType.get())
   with InsertionPoint(with_pdl.body):
     sequence = transform.SequenceOp(
-        transform.FailurePropagationMode.PROPAGATE, with_pdl.bodyTarget)
+        transform.FailurePropagationMode.PROPAGATE, [], with_pdl.bodyTarget)
     with InsertionPoint(sequence.body):
-      m1 = transform.PDLMatchOp(sequence.bodyTarget, "first")
-      m2 = transform.PDLMatchOp(sequence.bodyTarget, "second")
+      m1 = transform.PDLMatchOp(pdl.OperationType.get(), sequence.bodyTarget, "first")
+      m2 = transform.PDLMatchOp(pdl.OperationType.get(), sequence.bodyTarget, "second")
       transform.ReplicateOp(m1, [m2])
       transform.YieldOp()
   # CHECK-LABEL: TEST: testReplicateOp
   # CHECK: %[[FIRST:.+]] = pdl_match
   # CHECK: %[[SECOND:.+]] = pdl_match
   # CHECK: %{{.*}} = replicate num(%[[FIRST]]) %[[SECOND]]

mlir/test/python/dialects/transform_loop_ext.py

with InsertionPoint(module.body):
       print("\nTEST:", f.__name__)
       f()
     print(module)
   return f
 
 
 @run
 def getParentLoop():
-  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE)
+  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE, [], pdl.OperationType.get())
   with InsertionPoint(sequence.body):
     loop.GetParentForOp(sequence.bodyTarget, num_loops=2)
     transform.YieldOp()
   # CHECK-LABEL: TEST: getParentLoop
   # CHECK: = transform.loop.get_parent_for %
   # CHECK: num_loops = 2
 
 
 @run
 def loopOutline():
-  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE)
+  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE, [], pdl.OperationType.get())
   with InsertionPoint(sequence.body):
     loop.LoopOutlineOp(sequence.bodyTarget, func_name="foo")
     transform.YieldOp()
   # CHECK-LABEL: TEST: loopOutline
   # CHECK: = transform.loop.outline %
   # CHECK: func_name = "foo"
 
 
 @run
 def loopPeel():
-  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE)
+  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE, [], pdl.OperationType.get())
   with InsertionPoint(sequence.body):
     loop.LoopPeelOp(sequence.bodyTarget)
     transform.YieldOp()
   # CHECK-LABEL: TEST: loopPeel
   # CHECK: = transform.loop.peel %
 
 
 @run
 def loopPipeline():
-  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE)
+  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE, [], pdl.OperationType.get())
   with InsertionPoint(sequence.body):
     loop.LoopPipelineOp(sequence.bodyTarget, iteration_interval=3)
     transform.YieldOp()
   # CHECK-LABEL: TEST: loopPipeline
   # CHECK: = transform.loop.pipeline %
   # CHECK-DAG: iteration_interval = 3
   # CHECK-DAG: read_latency = 10
 
 
 @run
 def loopUnroll():
-  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE)
+  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE, [], pdl.OperationType.get())
   with InsertionPoint(sequence.body):
     loop.LoopUnrollOp(sequence.bodyTarget, factor=42)
     transform.YieldOp()
   # CHECK-LABEL: TEST: loopUnroll
   # CHECK: transform.loop.unroll %
   # CHECK: factor = 42

mlir/test/python/dialects/transform_structured_ext.py

with InsertionPoint(module.body):
       print("\nTEST:", f.__name__)
       f()
     print(module)
   return f
 
 
 @run
 def testDecompose():
-  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE)
+  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE, [], pdl.OperationType.get())
   with InsertionPoint(sequence.body):
     structured.DecomposeOp(sequence.bodyTarget)
     transform.YieldOp()
   # CHECK-LABEL: TEST: testDecompose
   # CHECK: transform.sequence
   # CHECK: transform.structured.decompose
 
 
 @run
 def testGeneralize():
-  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE)
+  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE, [], pdl.OperationType.get())
   with InsertionPoint(sequence.body):
     structured.GeneralizeOp(sequence.bodyTarget)
     transform.YieldOp()
   # CHECK-LABEL: TEST: testGeneralize
   # CHECK: transform.sequence
   # CHECK: transform.structured.generalize
 
 
 @run
 def testInterchange():
-  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE)
+  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE, [], pdl.OperationType.get())
   with InsertionPoint(sequence.body):
     structured.InterchangeOp(
         sequence.bodyTarget,
         iterator_interchange=[
             IntegerAttr.get(IntegerType.get_signless(64), 1), 0
         ])
     transform.YieldOp()
   # CHECK-LABEL: TEST: testInterchange
   # CHECK: transform.sequence
   # CHECK: transform.structured.interchange
   # CHECK: iterator_interchange = [1, 0]
 
 
 @run
 def testMultitileSizes():
-  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE)
+  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE, [], pdl.OperationType.get())
   with InsertionPoint(sequence.body):
     structured.MultiTileSizesOp(
         sequence.bodyTarget, dimension=1, target_size=42)
     transform.YieldOp()
   # CHECK-LABEL: TEST: testMultitileSizes
   # CHECK: transform.sequence
   # CHECK: transform.structured.multitile_sizes
   # CHECK-DAG: dimension = 1
   # CHECK-DAG: target_size = 42
 
 
 @run
 def testPad():
-  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE)
+  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE, [], pdl.OperationType.get())
   with InsertionPoint(sequence.body):
     structured.PadOp(
         sequence.bodyTarget,
         padding_values=[FloatAttr.get_f32(42.0)],
         padding_dimensions=[1],
         transpose_paddings=[[1, 0]])
     transform.YieldOp()
   # CHECK-LABEL: TEST: testPad
   # CHECK: transform.sequence
   # CHECK: transform.structured.pad
   # CHECK-DAG: padding_values = [4.200000e+01 : f32]
   # CHECK-DAG: padding_dimensions = [1]
   # CHECK-DAG: transpose_paddings = {{\[}}[1, 0]]
   # CHECK-DAG: hoist_paddings = []
   # CHECK-DAG: pack_paddings = []
 
 
 @run
 def testScalarize():
-  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE)
+  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE, [], pdl.OperationType.get())
   with InsertionPoint(sequence.body):
     structured.ScalarizeOp(sequence.bodyTarget)
     transform.YieldOp()
   # CHECK-LABEL: TEST: testScalarize
   # CHECK: transform.structured.scalarize
 
 
 @run
 def testSplit():
-  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE)
+  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE, [], pdl.OperationType.get())
   with InsertionPoint(sequence.body):
     split = structured.SplitOp(sequence.bodyTarget, dimension=1, split_point=42)
     structured.SplitOp(
         split.results[0], dimension=3, split_point=split.results[1])
     transform.YieldOp()
   # CHECK-LABEL: TEST: testSplit
   # CHECK: %[[F:.+]], %[[S:.+]] = transform.structured.split %{{.*}} after 42 {dimension = 1
   # CHECK: transform.structured.split %[[F]] after %[[S]] {dimension = 3
 
 
 @run
 def testTileCompact():
-  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE)
+  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE, [], pdl.OperationType.get())
   with InsertionPoint(sequence.body):
     structured.TileOp(sequence.bodyTarget, sizes=[4, 8], interchange=[0, 1])
     transform.YieldOp()
   # CHECK-LABEL: TEST: testTileCompact
   # CHECK: transform.sequence
   # CHECK: %{{.+}}, %{{.+}}:2 = transform.structured.tile %{{.*}}[4, 8]
   # CHECK: interchange = [0, 1]
 
 
 @run
 def testTileAttributes():
-  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE)
+  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE, [], pdl.OperationType.get())
   attr = ArrayAttr.get(
       [IntegerAttr.get(IntegerType.get_signless(64), x) for x in [4, 8]])
   ichange = ArrayAttr.get(
       [IntegerAttr.get(IntegerType.get_signless(64), x) for x in [0, 1]])
   with InsertionPoint(sequence.body):
     structured.TileOp(sequence.bodyTarget, sizes=attr, interchange=ichange)
     transform.YieldOp()
   # CHECK-LABEL: TEST: testTileAttributes
   # CHECK: transform.sequence
   # CHECK: %{{.+}}, %{{.+}}:2 = transform.structured.tile %{{.*}}[4, 8]
   # CHECK: interchange = [0, 1]
 
 
 @run
 def testTileZero():
-  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE)
+  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE, [], pdl.OperationType.get())
   with InsertionPoint(sequence.body):
     structured.TileOp(
         sequence.bodyTarget, sizes=[4, 0, 2, 0], interchange=[0, 1, 2, 3])
     transform.YieldOp()
   # CHECK-LABEL: TEST: testTileZero
   # CHECK: transform.sequence
   # CHECK: %{{.+}}, %{{.+}}:2 = transform.structured.tile %{{.*}}[4, 0, 2, 0]
   # CHECK: interchange = [0, 1, 2, 3]
 
 
 @run
 def testTileDynamic():
-  with_pdl = transform.WithPDLPatternsOp()
+  with_pdl = transform.WithPDLPatternsOp(pdl.OperationType.get())
   with InsertionPoint(with_pdl.body):
-    sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE,
+    sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE, [],
                                     with_pdl.bodyTarget)
     with InsertionPoint(sequence.body):
-      m1 = transform.PDLMatchOp(sequence.bodyTarget, "first")
-      m2 = transform.PDLMatchOp(sequence.bodyTarget, "second")
+      m1 = transform.PDLMatchOp(pdl.OperationType.get(), sequence.bodyTarget, "first")
+      m2 = transform.PDLMatchOp(pdl.OperationType.get(), sequence.bodyTarget, "second")
       structured.TileOp(sequence.bodyTarget, sizes=[m1, 3, m2, 0])
       transform.YieldOp()
   # CHECK-LABEL: TEST: testTileDynamic
   # CHECK: %[[FIRST:.+]] = pdl_match
   # CHECK: %[[SECOND:.+]] = pdl_match
   # CHECK: %{{.+}}, %{{.+}}:3 = transform.structured.tile %{{.*}}[%[[FIRST]], 3, %[[SECOND]], 0]
 
 
 @run
 def testVectorize():
-  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE)
+  sequence = transform.SequenceOp(transform.FailurePropagationMode.PROPAGATE, [], pdl.OperationType.get())
   with InsertionPoint(sequence.body):
     structured.VectorizeOp(sequence.bodyTarget, vectorize_padding=True)
     transform.YieldOp()
   # CHECK-LABEL: TEST: testVectorize
   # CHECK: transform.sequence
   # CHECK: = transform.structured.vectorize
   # CHECK: {vectorize_padding}

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

 )
 
 td_library(
     name = "TransformDialectTdFiles",
     srcs = glob(["include/mlir/Dialect/Transform/IR/*.td"]),
     deps = [
         ":CastInterfacesTdFiles",
         ":ControlFlowInterfacesTdFiles",
+        ":InferTypeOpInterfaceTdFiles",
         ":OpBaseTdFiles",
         ":PDLDialectTdFiles",
         ":SideEffectInterfacesTdFiles",
     ],
 )
 
 gentbl_cc_library(
     name = "TransformDialectEnumsIncGen",