Transform Dialect

Fine-grain transformation control dialect.

Disclaimer ¶

This dialect is actively developed and may change frequently.

To decrease the maintenance burden and churn, please post a description of the intended use case on the MLIR forum. A few in-tree use cases are currently supported:

• high-level transformations on “structured ops” (i.e. ops that operate on chunks of data in a way that can be decomposed into operations on smaller chunks of data and control flow) in Linalg, Tensor and Vector dialects;
• loop transformations in the SCF dialect.

Overview ¶

This dialect provides operations that can be used to control transformation of the IR using a different portion of the IR. It refers to the IR being transformed as payload IR, and to the IR guiding the transformation as transform IR.

The main use case for this dialect is orchestrating fine-grain transformations on individual operations or sets thereof. For example, it may involve finding loop-like operations with specific properties (e.g., large size) in the payload IR, applying loop tiling to those and only those operations, and then applying loop unrolling to the inner loops produced by the previous transformations. As such, it is not intended as a replacement for the pass infrastructure, nor for the pattern rewriting infrastructure. In the most common case, the transform IR will be processed and applied to the payload IR by a pass. Transformations expressed by the transform dialect may be implemented using the pattern infrastructure or any other relevant MLIR component.

The following IR gives a rough idea of what the operations in this dialect may look like:

%0 = transform.loop.find { size > 42 } : !transform.interface<tileable>
%1:2 = transform.loop.tile %0 { tile_sizes = [2,3,4] }
      : (!transform.interface<tileable>)
    -> (!transform.op<loop>, !transform.op<loop>)
transform.loop.unroll %1#1 : !transform.op<loop>

The values used in the Transform dialect, also referred to as handles, correspond to (groups of) operations in the payload IR. In the example above, %0 corresponds to the set of loops found in the payload IR that satisfy the condition, and %1 correspond to groups of outer and inner loops, respectively, produced by the tiling transformation.

A transform handle such as %0 may be associated with multiple payload operations. This is conceptually a set of operations and no assumptions should be made about the order of ops unless specified otherwise by the operation. Most Transform IR ops support operand values that are mapped to multiple operations. They usually apply the respective transformation for every mapped op (“batched execution”). Deviations from this convention are described in the documentation of Transform IR ops.

The handle values have transform IR types. These types describe properties of payload IR operations associated with the value that are known to the transform dialect, for example, all associated payload operations implement a “TileableOp” interface, or have a specific “loop” kind. These properties are used to statically indicate pre- and post-conditions of a transformation connected to a Transform dialect operation. The conditions are verified when payload IR operations are first associated with a transform handle. By convention, Transform dialect operations are expected to indicate narrow preconditions for their operands by enforcing operand type constraints in the their definitions and verifiers. On the contrary, operations are expected to have few constraints on their results. Specific instances of a transform operation can then be created with a more restricted result type than the constraint in the operation (e.g., the “find” operation only constrains the result type to be a transform IR type while its concrete instance can have a type with stricter constraints such as implementing the “tilable” interface). The verification will then happen at transform execution time. This approach allows one to capture payload IR operation properties in the transform IR without resorting to excessive use of type casts or coupling dialect extensions between themselves. It is a trade-off between verbosity/complexity and static hardening, which can be revised in the future.

Overall, Transform IR ops are expected to be contained in a single top-level op. Such top-level ops specify how to apply the transformations described by the operations they contain, e.g., transform.sequence executes transformations one by one and fails if any of them fails. Such ops are expected to have the PossibleTopLevelTransformOpTrait and may be used without arguments.

A program transformation expressed using the Transform dialect can be programmatically triggered by calling:

LogicalResult transform::applyTransforms(Operation *payloadRoot,
                                          TransformOpInterface transform,
                                          const TransformOptions &options);

that applies the transformations specified by the top-level transform to payload IR contained in payloadRoot.

Dialect Extension Mechanism ¶

This dialect is designed to be extensible, that is, clients of this dialect are allowed to inject additional operations into this dialect using the TransformDialectExtension mechanism. This allows the dialect to avoid a dependency on the implementation of the transformation as well as to avoid introducing dialect-specific transform dialects. In the example above, the operations may have been injected by a notional loop dialect rather than defined in this dialect, hence the common prefix.

It is recommended to prefix injected operations with one or several dot-separated words that indicate which extension adds them. For dialect-specific transformations, the prefix is naturally the name of the dialect, e.g., transform.affine.reschedule. For dialect-agnostic transformations (typically implemented using interfaces), the prefix may be derived from the interface name or from a common concept, e.g., transform.loop.tile may apply to any loop-like operation that implements TileableOpInterface. The C++ classes for the dialect extension should include the prefix in their name, e.g., AffineTransformDialectExtension or LoopTransformDialectExtension in the cases above. Unprefixed operation names are reserved for ops defined directly in the Transform dialect.

Operations injected into the dialect must:

• Implement the TransformOpInterface to execute the corresponding transformation on the payload IR.

• Implement the MemoryEffectsOpInterface to annotate the effects of the transform IR operation on the payload IR as well as on the mapping between transform IR values and payload IR operations. See below for the description of available effects.

The presence of interface implementations is checked at runtime when the dialect is loaded to allow for those implementations to be supplied by separate dialect extensions if desired.

Side Effects ¶

The Transform dialect relies on MLIR side effect modelling to enable optimization of the transform IR. More specifically, it provides several side effect resource objects and expects operations to describe their effects on these resources.

• TransformMappingResource - side effect resource corresponding to the mapping between transform IR values and payload IR operations.

  • An Allocate effect from this resource means creating a new mapping entry, it is always accompanied by a Write effect.

  • A Read effect from this resource means accessing the mapping.

  • A Free effect on this resource indicates the removal of the mapping entry, typically after a transformation that modifies the payload IR operations associated with one of the transform IR operation’s operands. It is always accompanied by a Read effect.

• PayloadIRResource - side effect resource corresponding to the payload IR itself.

  • A Read effect from this resource means accessing the payload IR.

  • A Write effect on this resource means mutating the payload IR. It is almost always accompanied by a Read.

The typical flow of values in the transform IR is as follows. Most operations produce new transform IR values and immediately associate them with a list of payload IR operations. This corresponds to Allocate and Write effects on the TransformMappingResource, and often requires at least a Read effect on the PayloadIRResource. Transform operations that only inspect the payload IR to produce new handles are usually limited to these effects on their operands. Transform operations that mutate the payload IR are thought to consume the handles provided as operands, that is have the Read and Free effects on them. As with the usual memory effects, using a value after it was freed is incorrect. In case of the transform IR, this value is likely associated with payload IR operations that were modified or even removed by the transformation, so it is meaningless to refer to them. When further transformations are desired, the transform operations can return new handles that can be read or consumed by subsequent operations.

Execution Model ¶

The transformation starts at the user-specified top-level transform IR operation and applies to some user-specified payload IR scope, identified by the payload IR op that contains the IR to transform. It is the responsibility of the user to properly select the scope and/or to avoid the transformations to modify the IR outside of the given scope. The top-level transform IR operation may contain further transform operations and execute them in the desired order.

Transformation application functions produce a tri-state status:

• success;
• recoverable (silenceable) failure;
• irrecoverable failure.

Transformation container operations may intercept recoverable failures and perform the required recovery steps thus succeeding themselves. On the other hand, they must propagate irrecoverable failures. For such failures, the diagnostics are emitted immediately whereas their emission is postponed for recoverable failures. Transformation container operations may also fail to recover from a theoretically recoverable failure, in which case they can either propagate it to their parent or emit the diagnostic and turn the failure into an irrecoverable one. A recoverable failure produced by applying the top-level transform IR operation is considered irrecoverable.

Transformation container operations are allowed to “step over” some nested operations if the application of some previous operation produced a failure. This can be conceptually thought of as having a global “recoverable error register” that is read/write accessed by each transform operation as a side effect. The transformation is skipped if the register already contains an error description, and the control flow proceeds to the following operation.

Note that a silenceable failure, if emitted, is a compiler error rather than a warning. Transformations are expected to produce silenceable failures if they haven’t yet modified the payload IR, i.e. when reporting a precondition failure, and an irrecoverable failure when they modified the IR in a way that is contrary to the semantics of the transform operation or would fail a postcondition. Some “navigation” operations that identify payload IR targets for the following transformation may have a conceptual “failure to match” that is considered a successful execution in the execution model but results in handles associated with empty payload IR operation lists.

Handle Invalidation ¶

The execution model of the transform dialect allows a payload IR operation to be associated with multiple handles as well as nested payload IR operations to be associated with different handles. A transform IR operation that consumes a handle automatically invalidates all the other handles associated with the same payload IR operations, or with any of their descendants, as the consumed handle. Note that the entire handle is invalidated, even if some of the payload IR operations associated with it or their ancestors were not associated with the consumed handle. Any use of the invalidated handle results in undefined behavior since the payload IR operations associated with it are likely to have been mutated or erased. The mere fact of the handle being invalidated does not trigger undefined behavior, only its appearance as an operand does.

The Transform dialect infrastructure has the capability of checking whether the transform IR op operand is invalidated before applying the transformation. However, such a check is computationally expensive and must be enabled explicitly through TransformOptions. Additionally, the transform-dialect-check-uses pass emits warnings when a handle may be used after it has been consumed, but does so abstractly, without processing the payload IR.

Intended Use and Integrations ¶

The transformation control infrastructure provided by this dialect is positioned roughly between rewrite patterns and passes. A transformation that is executed by a transform operation is likely to be sufficiently complex to require at least a set of patterns to be implemented. It is also expected to be more focused than a pass: a pass typically applies identical transformations everywhere in the IR, a transform dialect-controlled transformation would apply to a small subset of operations selected, e.g., by a pattern-matching operation or generated by a previous transformation. It is discouraged, although technically possible, to run a pass pipeline as part of the transform op implementation.

One of the main scenarios for using this dialect is fine-grain chaining of transformations. For example, a loop-like operation may see its iteration domain split into two parts, implemented as separate loops (transformation known as index-set splitting), each of which is then transformed differently (e.g., the first loop is tiled and the second unrolled) with the necessary enabling and cleanup patterns around the main transformation:

// <generate %loop, e.g., by pattern-matching>
// ...
%parts:2 = transform.loop.split %loop { upper_bound_divisible_by = 8 }
transform.loop.tile %parts#0 { tile_sizes = [8] }
transform.loop.unroll %parts#1 { full }

This composition would have been difficult to implement as separate passes since the hypothetical “tiling” and “unrolling” pass would need to somehow differentiate between the parts of the loop produced by the previous pass (both are the same operation, and it is likely undesirable to pollute the operation with pass-specific information). Implementing passes that run the combined transformation would have run into the combinatorial explosion issue due to multiple possible transform compositions or into the need for deep pass parameterization, the ultimate form of which is an ad-hoc dialect to specify which transformations the pass should run. The transform dialect provides a uniform, extensible mechanism for controlling transformations in such cases.

The transform dialect is supposed to be consumed by an “interpreter” pass that drives the application of transformations. To ensure extensibility and composability, this pass is not expected to actually perform the transformations specified by the ops. Instead, the transformations are implemented by the transform ops themselves via TransformOpInterface. The pass serves as the entry point, handles the flow of transform operations and takes care of bookkeeping. As such, the transform dialect does not provide the interpreter pass. Instead, it provides a set of utilities that can be used by clients to define their own interpreter passes or as part of a more complex pass. For example, the mapping between values in the transform IR and operations in the payload IR, or the function that applies the transformations specified by ops in the given block sequentially. Note that a transform op may have regions with further transform ops in them, with the op itself guiding how to dispatch the transformation control flow to those regions. This approach allows clients to decide on the relative location of the transform IR in their input (e.g., nested modules, separate modules, optional regions to certain operations, etc.), register additional transform operations and perform client-specific bookkeeping.

Effects on the Infrastructure ¶

Although scoped to a single dialect, this functionality conceptually belongs to the MLIR infrastructure. It aims to be minimally intrusive and opt-in.

Some infrastructural components may grow extra functionality to support the transform dialect. In particular, the pattern infrastructure may add extra hooks to identify the “main results” of a transformation or to notify external observers about changes made to certain operations. These are not expected to affect the existing uses of the infrastructure.

For the sake of reusability, transformations should be implemented as utility functions that are called from the interface methods of transform ops rather than having the methods directly act on the payload IR.

Type Definitions ¶

AnyOpType ¶

Syntax: !transform.any_op

Transform IR handle that can be associated with a list of arbitrary Payload IR operations.

OperationType ¶

Syntax:

!transform.op<
  ::llvm::StringRef   # operation_name
>

Transform IR handle that can be associated with a list of Payload IR operations with the specified operation name.

Parameters: ¶

Core Operations ¶

transform.alternatives (::mlir::transform::AlternativesOp) ¶

Attempts sequences of transforms until one succeeds

Syntax:

operation ::= `transform.alternatives` ($scope^ `:` type($scope))? (`->` type($results)^)? attr-dict-with-keyword regions

This op may have an arbitrary number of regions, each of which represents a sequence of transform operations to be applied to the same payload IR. The regions are visited in order of appearance, and transforms in them are applied in their respective order of appearance. If one of these transforms fails to apply, the remaining ops in the same region are skipped an the next region is attempted. If all transformations in a region succeed, the remaining regions are skipped and the entire “alternatives” transformation succeeds. If all regions contained a failing transformation, the entire “alternatives” transformation fails.

It is up to the nested operations to define which errors are “recoverable” (or “silenceable”) and allow another alternatives to be attempted, and which errors should be propagated without attempting the other alternatives.

The single operand of this operation is the scope in which the alternative transformation sequences are attempted, that is, an operation in the payload IR that contains all the other operations that may be modified by the transformations. The scope operation must be isolated from above. There is no check that the transforms are indeed scoped as their “apply” methods can be arbitrarily complex. Therefore it is the responsibility of the user to ensure that the transforms are scoped correctly, or to produce an irrecoverable error and thus abort the execution without attempting the remaining alternatives. Note that the payload IR outside of the given scope is not necessarily in the valid state, or even accessible to the transformation.

The changes to the IR within the scope performed by transforms in the failed alternative region are reverted before attempting the next region. Practically, this is achieved by cloning the scope. Therefore it is advised to limit the scope as much as possible and place the most likely alternatives early in the region list. The operation is also isolated from above and requires rediscovering the operations within the given scope to avoid additional handle invalidation. The latter restriction may be lifted in the future.

Each of the regions may yield transform IR handles. The handles of the first successful alternative region are returned as the results of the “alternatives” op. Therefore, each alternative region must yield the same number of results, which should also match the number and the types of the “alternatives” op results.

Remark: this op allows one to implement a simple “try” construct as follows:

%result = transform.alternatives %scope {
^bb0(%arg0: !pdl.operation):
  // Try a fallible transformation.
  %0 = transform.fallible %arg0 // ...
  // If succeeded, yield the the result of the transformation.
  transform.yield %0 : !pdl.operation
}, {
^bb0(%arg0: !pdl.operation):
  // Otherwise, the second alternative is tried and it always succeeds by
  // returning the original handle.
  transform.yield %arg0 : !pdl.operation
}

Traits: FunctionalStyleTransformOpTrait, IsolatedFromAbove, PossibleTopLevelTransformOpTrait, SingleBlockImplicitTerminator<::mlir::transform::YieldOp>

Interfaces: MemoryEffectsOpInterface, RegionBranchOpInterface, TransformOpInterface

Operands: ¶

Results: ¶

transform.cast (::mlir::transform::CastOp) ¶

Syntax:

operation ::= `transform.cast` $input attr-dict `:` type($input) `to` type($output)

Traits: TransformEachOpTrait

Interfaces: CastOpInterface, MemoryEffectOpInterface, TransformOpInterface

Operands: ¶

Results: ¶

transform.foreach (::mlir::transform::ForeachOp) ¶

Executes the body for each payload op

Syntax:

operation ::= `transform.foreach` $target `:` type($target) (`->` type($results)^)? $body attr-dict

This op has exactly one region with exactly one block (“body”). The body is executed for each payload op that is associated to the target operand in an unbatched fashion. I.e., the block argument (“iteration variable”) is always mapped to exactly one payload op.

This op always reads the target handle. Furthermore, it consumes the handle if there is a transform op in the body that consumes the iteration variable. This op does not return anything.

The transformations inside the body are applied in order of their appearance. During application, if any transformation in the sequence fails, the entire sequence fails immediately leaving the payload IR in potentially 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.

Traits: SingleBlockImplicitTerminator<::mlir::transform::YieldOp>

Interfaces: MemoryEffectOpInterface, RegionBranchOpInterface, TransformOpInterface

Operands: ¶

Results: ¶

transform.get_closest_isolated_parent (::mlir::transform::GetClosestIsolatedParentOp) ¶

Gets handles to the closest isolated-from-above parents

Syntax:

operation ::= `transform.get_closest_isolated_parent` $target attr-dict `:` functional-type(operands, results)

The handles defined by this Transform op correspond to the closest isolated from above ancestor of the Payload IR operations associated with its operand. If any of the given Payload IR ops has no such parent (unlikely as there usually is a top-level ModuleOp), the transformation is considered to have failed.

Ancestor ops follow the same order as the ops associated with the operand, except for potential duplicates (multiple Payload IR ops associated 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.

Traits: NavigationTransformOpTrait

Interfaces: MemoryEffectsOpInterface, TransformOpInterface

Operands: ¶

Results: ¶

transform.get_producer_of_operand (::mlir::transform::GetProducerOfOperand) ¶

Get handle to the producer of this operation’s operand number

Syntax:

operation ::= `transform.get_producer_of_operand` $target `[` $operand_number `]` attr-dict `:` functional-type(operands, results)

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.

Traits: NavigationTransformOpTrait

Interfaces: MemoryEffectsOpInterface, TransformOpInterface

Attributes: ¶

Operands: ¶

Results: ¶

transform.merge_handles (::mlir::transform::MergeHandlesOp) ¶

Merges handles into one pointing to the union of payload ops

Syntax:

operation ::= `transform.merge_handles` ($deduplicate^)? $handles attr-dict `:` type($result)

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.

Traits: SameOperandsAndResultType

Interfaces: MemoryEffectOpInterface, TransformOpInterface

Attributes: ¶

Operands: ¶

Results: ¶

transform.pdl_match (::mlir::transform::PDLMatchOp) ¶

Finds ops that match the named PDL pattern

Syntax:

operation ::= `transform.pdl_match` $pattern_name `in` $root attr-dict `:` functional-type(operands, results)

Find Payload IR ops nested within the Payload IR op associated with the operand that match the PDL pattern identified by its name. The pattern is expected to be defined in the closest surrounding WithPDLPatternsOp.

Produces a Transform IR value associated with the list of Payload IR ops that matched the pattern. The order of results in the list is that of the Operation::walk, clients are advised not to rely on a specific order though. If the operand is associated with multiple Payload IR ops, finds matching ops nested within each of those and produces a single list containing all 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.

Interfaces: MemoryEffectOpInterface, TransformOpInterface

Attributes: ¶

Operands: ¶

Results: ¶

transform.print (::mlir::transform::PrintOp) ¶

Dump each payload op

Syntax:

operation ::= `transform.print` $target attr-dict (`:` type($target)^)?

This op dumps each payload op that is associated with the target operand to stderr. It also prints the name string attribute. If no target is specified, the top-level op is dumped.

This op is useful for printf-style debugging.

Interfaces: MemoryEffectOpInterface, TransformOpInterface

Attributes: ¶

Operands: ¶

transform.replicate (::mlir::transform::ReplicateOp) ¶

Lists payload ops multiple times in the new handle

Syntax:

operation ::= `transform.replicate` `num` `(` $pattern `)` $handles attr-dict `:` type($pattern) `,` type($handles)

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].

This transformation is useful to “align” the sizes of payload IR lists before a transformation that expects, e.g., identically-sized lists. For example, a transformation may be parameterized by same notional per-target size computed at runtime and supplied as another handle, the replication allows this size to be computed only once and used for every target instead of replicating the computation itself.

Note that it is undesirable to pass a handle with duplicate operations to 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.

Interfaces: MemoryEffectOpInterface, TransformOpInterface

Operands: ¶

Results: ¶

transform.sequence (::mlir::transform::SequenceOp) ¶

Contains a sequence of other transform ops to apply

Syntax:

operation ::= `transform.sequence` ($root^ `:` type($root))? (`->` type($results)^)? `failures` `(` $failure_propagation_mode `)` attr-dict-with-keyword regions

The transformations indicated by the sequence are applied in order of their appearance. Each value produced by a transformation within the sequence corresponds to an operation or a group of operations in the payload IR. The behavior of the operation when a nested transformation produces a silenceable error is controlled by the failure_propagation_mode attribute. When set to propagate, the failure of any nested transformation in the sequence implies immediate failure of the entire sequence with a silenceable error, and no further transformation is attempted. When set to suppress, silenceable errors in nested operations are ignored and further transformations are applied. Beware that even silenceable errors may leave the payload IR in a state unsuitable for further transformations. It is the responsibility of the caller to ensure the following transformations are robust enough when errors are suppressed. Definite errors reported by nested transformations abort the sequence regardless of the propagation mode. The set of modes may be extended in the future, e.g., to collect silenceable errors and report them after attempting all transformations in the sequence.

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.

The body of the sequence terminates with an implicit or explicit transform.yield op. The operands of the terminator are returned as the results of the sequence op.

Traits: PossibleTopLevelTransformOpTrait, SingleBlockImplicitTerminator<::mlir::transform::YieldOp>

Interfaces: MemoryEffectOpInterface, OpAsmOpInterface, RegionBranchOpInterface, TransformOpInterface

Attributes: ¶

Operands: ¶

Results: ¶

transform.split_handles (::mlir::transform::SplitHandlesOp) ¶

Splits handles from a union of payload ops to a list

Syntax:

operation ::= `transform.split_handles` $handle `in` `[` $num_result_handles `]`
              attr-dict `:` functional-type(operands, results)

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.

Traits: FunctionalStyleTransformOpTrait

Interfaces: MemoryEffectOpInterface, TransformOpInterface

Attributes: ¶

Operands: ¶

Results: ¶

transform.with_pdl_patterns (::mlir::transform::WithPDLPatternsOp) ¶

Contains PDL patterns available for use in transforms

Syntax:

operation ::= `transform.with_pdl_patterns` ($root^ `:` type($root))? attr-dict-with-keyword regions

This op contains a set of named PDL patterns that are available for the Transform dialect operations to be used for pattern matching. For example, PDLMatchOp can be used to produce a Transform IR value associated with all Payload IR operations that match the pattern as follows:

transform.with_pdl_patterns {
^bb0(%arg0: !pdl.operation):
  pdl.pattern @my_pattern : benefit(1) {
    %0 = pdl.operation //...
    // Regular PDL goes here.
    pdl.rewrite %0 with "transform.dialect"
  }

  sequence %arg0 failures(propagate) {
  ^bb0(%arg1: !pdl.operation):
    %1 = pdl_match @my_pattern in %arg1
    // Use %1 as handle
  }
}

Note that the pattern is expected to finish with a pdl.rewrite terminator that points to the custom rewriter named “transform.dialect”. The rewriter actually does nothing, but the transform application will keep track of the operations that matched the pattern.

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.

Traits: NoTerminator, PossibleTopLevelTransformOpTrait, RecursiveMemoryEffects, SymbolTable

Interfaces: OpAsmOpInterface, TransformOpInterface

Operands: ¶

transform.yield (::mlir::transform::YieldOp) ¶

Yields operation handles from a transform IR region

Syntax:

operation ::= `transform.yield` operands attr-dict (`:` type($operands)^)?

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.

Traits: Terminator

Operands: ¶

Bufferization Transform Operations ¶

transform.bufferization.one_shot_bufferize (::mlir::transform::OneShotBufferizeOp) ¶

Syntax:

operation ::= `transform.bufferization.one_shot_bufferize` (`layout` `{` $function_boundary_type_conversion^ `}`)?
              $target attr-dict

Indicates that the given target op should be bufferized with One-Shot Bufferize. The bufferization can be configured with various attributes that corresponding to options in BufferizationOptions and the one-shot-bufferize pass. More information can be found in the pass documentation.

If target_is_module is set, target must be a module. In that case the target handle can be reused by other transform ops. When bufferizing other ops, the target handled is freed after bufferization and can no longer be used.

Note: Only ops that implement BufferizableOpInterface are bufferized. All other ops are ignored if allow_unknown_ops. If allow_unknown_ops is unset, this transform fails when an unknown/non-bufferizable op is found. Many ops implement BufferizableOpInterface via an external model. These external models must be registered when applying this transform op; otherwise, said ops would be considered non-bufferizable.

Interfaces: MemoryEffectOpInterface, TransformOpInterface

Attributes: ¶

Operands: ¶

GPU Transform Operations ¶

transform.gpu.map_foreach_to_blocks (::mlir::transform::MapForeachToBlocks) ¶

Syntax:

operation ::= `transform.gpu.map_foreach_to_blocks` $target attr-dict

Target the gpu_launch op and rewrite the top level scf.foreach_thread to distributed gpu.block_id attribute. If generate_gpu_launch attribute is set, then first generates gpu_launch and moves the top level scf.foreach_thread inside.

The operation searches top level scf.foreach_thread ops under gpu_launch and maps each such op to GPU blocks. Mapping is one-to-one and the induction variables of scf.foreach_thread are rewritten to gpu.block_id according to the thread_dim_apping attribute.

Dynamic, scf.foreach_thread trip counts are currently not supported. Dynamic block dim sizes are currently not supported.

Only bufferized scf.foreach_thread are currently supported. Only scf.foreach_thread distributed to at most 3 dimensions are currently supported.

The operation alters the block size of the given gpu_launch using gridDim argument.

Return modes: ¶

This operation ignores non-gpu_launch ops and drops them in the return.

If any scf.foreach_thread with tensors is found, the transform definitely fails.

If all the scf.foreach_thread operations contained within the LaunchOp referred to by the target PDLOperation lower to GPU properly, the transform succeeds. Otherwise the transform definitely fails.

The returned handle points to the same LaunchOp operand, consuming it and producing a new SSA value to satisfy chaining and linearity of the IR properties.

Traits: FunctionalStyleTransformOpTrait, TransformEachOpTrait

Interfaces: MemoryEffectsOpInterface, TransformOpInterface

Attributes: ¶

Operands: ¶

Results: ¶

transform.gpu.map_nested_foreach_to_threads (::mlir::transform::MapNestedForeachToThreads) ¶

Syntax:

operation ::= `transform.gpu.map_nested_foreach_to_threads` $target attr-dict

Target the gpu.launch op and rewrite all scf.foreach_thread nested in it to distributed gpu.thread_id attribute.

The operation searches for scf.foreach_thread ops nested under target and maps each such op to GPU threads. Mapping is one-to-one and the induction variables of scf.foreach_thread are rewritten to gpu.thread_id according to the mapping attribute.

Sibling scf.foreach_thread are supported in which case, the union of the number of threads is computed and may result in predication.

Multiple scf.foreach_thread are supported per gpu.launch in which case, the max of all the threads is computed and taken for the global gpu.thread_id. If necessary, scf.foreach_thread that do not use the whole thread range result in predicated computations.

Dynamic scf.foreach_thread trip counts are currently not supported. Dynamic block dim sizes are currently not supported.

Only bufferized scf.foreach_thread are currently supported. Only scf.foreach_thread distributed to at most 3 dimensions are currently supported.

Barriers are inserted after each scf.foreach_thread op for now.

The operation alters the block size of the given gpu_launch using blockDim argument.

Return modes: ¶

This operation ignores non-gpu_launch ops and drops them in the return.

If any scf.foreach_thread with tensors is found, the transform definitely fails.

If all the scf.foreach_thread operations contained within the LaunchOp referred to by the target PDLOperation lower to GPU properly, the transform succeeds. Otherwise the transform definitely fails.

The returned handle points to the same LaunchOp operand, consuming it and producing a new SSA value to satisfy chaining and linearity of the IR properties.

Example: ¶

gpu.launch blocks(%bx, %by, %bz) in (%x = %0, %y = %1, %z = %2)
           threads(%tx, %ty, %tz) in (%tx = %3, %ty = %4, %tz = %5) {
  scf.foreach_thread (%i, %j) in (7, 9) {
    ... // body 1
  } {mapping = [#gpu.thread<x>, #gpu.thread<y>, #gpu.thread<z>]}
  scf.foreach_thread (%i) in (12) {
    ... // body 2
  } {mapping = [#gpu.thread<x>]}
  gpu.terminator
}

is translated to:

%bdimX = arith.constant 12 : index
%bdimY = arith.constant 9 : index
gpu.launch blocks(%bx, %by, %bz) in (%x = %0, %y = %1, %z = %2)
       threads(%tx, %ty, %tz) in (%tx = %bdimX, %ty = %bdimY, %tz = %5) {
  if (threadIdx.x < 9 && threadIdx.y < 7) {
    ... // body 1
  }
  gpu.barrier
  if (threadIdx.y < 1) {
    ... // body 2
  }
  gpu.barrier
  gpu.terminator
}

Traits: FunctionalStyleTransformOpTrait, TransformEachOpTrait

Interfaces: MemoryEffectsOpInterface, TransformOpInterface

Attributes: ¶

Operands: ¶

Results: ¶

Loop (SCF) Transform Operations ¶

transform.loop.get_parent_for (::mlir::transform::GetParentForOp) ¶

Gets a handle to the parent ‘for’ loop of the given operation

Syntax:

operation ::= `transform.loop.get_parent_for` $target attr-dict `:` functional-type(operands, results)

Produces a handle to the n-th (default 1) parent scf.for or affine.for (when the affine flag is true) loop for each Payload IR operation associated with the operand. Fails if such a loop cannot be found. The list of operations associated with the handle contains parent operations in the same order as the list associated with the operand, except for operations that are parents to more than one input which are only present once.

Traits: NavigationTransformOpTrait

Interfaces: MemoryEffectsOpInterface, TransformOpInterface

Attributes: ¶

Operands: ¶

Results: ¶

transform.loop.outline (::mlir::transform::LoopOutlineOp) ¶

Outlines a loop into a named function

Syntax:

operation ::= `transform.loop.outline` $target attr-dict `:` functional-type(operands, results)

Moves the loop into a separate function with the specified name and replaces the loop in the Payload IR with a call to that function. Takes care of forwarding values that are used in the loop as function arguments. If the operand is associated with more than one loop, each loop will be outlined into a separate function. The provided name is used as a base for forming actual function names following SymbolTable auto-renaming scheme to avoid duplicate symbols. Expects that all ops in the Payload IR have a SymbolTable ancestor (typically true because of the top-level module). Returns the handle to the list of outlined functions in the same order as the operand handle.

Traits: FunctionalStyleTransformOpTrait

Interfaces: MemoryEffectsOpInterface, TransformOpInterface

Attributes: ¶

Operands: ¶

Results: ¶

transform.loop.peel (::mlir::transform::LoopPeelOp) ¶

Peels the last iteration of the loop

Syntax:

operation ::= `transform.loop.peel` $target attr-dict `:` functional-type(operands, results)

Updates the given loop so that its step evenly divides its range and puts the remaining iteration into a separate loop or a conditional.

In the absence of sufficient static information, this op may peel a loop, even if the step always divides the range evenly at runtime.

Return modes ¶

This operation ignores non-scf::ForOp ops and drops them in the return.

This operation always succeeds and returns the scf::ForOp with the postcondition: “the loop trip count is divisible by the step”. This operation may return the same unmodified loop handle when peeling did not modify the IR (i.e. the loop trip count was already divisible).

Note that even though the Payload IR modification may be performed in-place, this operation consumes the operand handle and produces a new one.

TODO: Return both the peeled loop and the remainder loop.

Traits: FunctionalStyleTransformOpTrait, TransformEachOpTrait

Interfaces: MemoryEffectsOpInterface, TransformOpInterface

Attributes: ¶

Operands: ¶

Results: ¶

transform.loop.pipeline (::mlir::transform::LoopPipelineOp) ¶

Applies software pipelining to the loop

Syntax:

operation ::= `transform.loop.pipeline` $target attr-dict `:` functional-type(operands, results)

Transforms the given loops one by one to achieve software pipelining for each of them. That is, performs some amount of reads from memory before the loop rather than inside the loop, the same amount of writes into memory after the loop, and updates each iteration to read the data for a following iteration rather than the current one.

The amount is specified by the attributes.

The values read and about to be stored are transferred as loop iteration arguments. Currently supports memref and vector transfer operations as memory reads/writes.

Return modes ¶

This operation ignores non-scf::For ops and drops them in the return. If all the operations referred to by the target PDLOperation pipeline properly, the transform succeeds. Otherwise the transform silently fails. The return handle points to only the subset of successfully produced pipelined loops, which can be empty.

Traits: FunctionalStyleTransformOpTrait, TransformEachOpTrait

Interfaces: MemoryEffectsOpInterface, TransformOpInterface

Attributes: ¶

Operands: ¶

Results: ¶

transform.loop.unroll (::mlir::transform::LoopUnrollOp) ¶

Unrolls the given loop with the given unroll factor

Syntax:

operation ::= `transform.loop.unroll` $target attr-dict `:` type($target)

Unrolls each loop associated with the given handle to have up to the given number of loop body copies per iteration. If the unroll factor is larger than the loop trip count, the latter is used as the unroll factor instead.

Return modes ¶

This operation ignores non-scf::For, non-affine::For ops and drops them in the return. If all the operations referred to by the target PDLOperation unroll properly, the transform succeeds. Otherwise the transform silently fails.

Does not return handles as the operation may result in the loop being removed after a full unrolling.

Traits: FunctionalStyleTransformOpTrait, TransformEachOpTrait

Interfaces: MemoryEffectsOpInterface, TransformOpInterface

Attributes: ¶

Operands: ¶

Structured (Linalg) Transform Operations ¶

transform.structured.decompose (::mlir::transform::DecomposeOp) ¶

Syntax:

operation ::= `transform.structured.decompose` $target attr-dict

Decomposes named complex operations, such as higher-dimensional (depthwise) convolutions, into combinations of lower-dimensional equivalents when possible.

Return modes ¶

This operation ignores non-Linalg ops and drops them in the return. If all the operations referred to by the target PDLOperation decompose properly, the transform succeeds. Otherwise the transform silently fails. The return handle points to only the subset of successfully produced computational operations, which can be empty.

Traits: FunctionalStyleTransformOpTrait, TransformEachOpTrait

Interfaces: MemoryEffectsOpInterface, TransformOpInterface

Operands: ¶

Results: ¶

transform.structured.fuse_into_containing_op (::mlir::transform::FuseIntoContainingOp) ¶

Fuse a producer into a containing operation.

Syntax:

operation ::= `transform.structured.fuse_into_containing_op` $producer_op `into` $containing_op attr-dict

Fuses the producer_op into the containing_op. Returns a handle to the fused ops.

The producer is typically a slice of a tileable op (i.e., implements TilingInterface). In that case, this transform computes the accessed producer slice inside of the containing op (“tile and fuse”). Otherwise, the entire producer is cloned inside the containing op (“clone and fuse”).

The containing op handle must be associated with exactly one payload op. The producer op handle may be associated with multiple payload ops. This transform fuses producers one-by-one, always picking an unspecified producer that has at least one use inside the containing op among the producers.

Note: If a producer has multiple uses inside the containing op, it is currently tiled and/or cloned multiple times into the containing op. TODO: Reuse already fused OpResults instead of tiling/cloning a second time when possible. Fuse producers according to a topological sorting to achieve the largest amount of reuse.

Return modes ¶

If at least one producer could not be fused, this operation fails silently. This is the case when tiling fails or when no producer op could be found among the remaining producers that has at least one use within the containing op. I.e., “producers” that are not consumed within the containing op are rejected by this operation.

This operation reads and frees the producer handle. This operation reads the containing op handle.

Interfaces: TransformOpInterface

Operands: ¶

Results: ¶

transform.structured.fuse (::mlir::transform::FuseOp) ¶

Tiles the operations pointed to by the target handle and fuses their producers greedily using the options provided as attributes.

Traits: FunctionalStyleTransformOpTrait

Interfaces: MemoryEffectsOpInterface, TransformOpInterface

Attributes: ¶

Operands: ¶

Results: ¶

transform.structured.generalize (::mlir::transform::GeneralizeOp) ¶

Syntax:

operation ::= `transform.structured.generalize` $target attr-dict

Transforms a named structured operation into the generic form with the explicit attached region.

Return modes ¶

This operation ignores non-Linalg ops and drops them in the return. If all the operations referred to by the target PDLOperation generalize properly, the transform succeeds. Otherwise the transform silently fails. The return handle points to only the subset of successfully produced equivalent generic operations, which can be empty or contain the original ops if they were already in generic form.

Traits: FunctionalStyleTransformOpTrait, TransformEachOpTrait

Interfaces: MemoryEffectsOpInterface, TransformOpInterface

Operands: ¶

Results: ¶

transform.structured.interchange (::mlir::transform::InterchangeOp) ¶

Syntax:

operation ::= `transform.structured.interchange` $target attr-dict

Interchanges the iterators of the operations pointed to by the target handle using the iterator interchange attribute.

Return modes ¶

This operation ignores non-linalg::Generic ops and drops them in the return. This operation fails if the interchange attribute is invalid. If all the operations referred to by the target PDLOperation interchange properly, the transform succeeds. If any interchange fails, the transform definitely fails. The return handle points to only the subset of successfully produced interchanged operations, which can be empty.

Traits: FunctionalStyleTransformOpTrait, TransformEachOpTrait

Interfaces: MemoryEffectsOpInterface, TransformOpInterface

Attributes: ¶

Operands: ¶

Results: ¶

transform.structured.match (::mlir::transform::MatchOp) ¶

Syntax:

operation ::= `transform.structured.match` (`ops` `{` $ops^ `}`)?
              (`interface` `{` $interface^ `}`)?
              (`attributes` $op_attrs^)?
              (`filter_result_type` `=` $filter_result_type^)?
              `in` $target attr-dict

Match op with the specified constraints, within the target op.

The following constraints are supported:

• interface: an optional MatchInterfaceEnum specifying an enum representation for an interface to target.
• ops: an optional StrArrayAttr specifying the concrete name of an op. Multiple names can be specified. Matched ops must have one of specified names.
• attribute: the matched op must have all specified attributes (with their specified values).
• filter_result_type: the matched op must return exactly this one type.

Note: Only ops that satisfy all specified constraints are matched.

TODO: Extend with regions to allow a limited form of constraints.

Return modes ¶

This op traverses the ops nested under target and returns the handles to all the operations that match the requirements.

This op fails if the target is not a handle to exactly one operation. Otherwise it succeeds.

This operation does not consume the target handle and produces new handles: it is a navigation op.

Traits: NavigationTransformOpTrait

Interfaces: MemoryEffectsOpInterface, TransformOpInterface

Attributes: ¶

Operands: ¶

Results: ¶

transform.structured.multitile_sizes (::mlir::transform::MultiTileSizesOp) ¶

Syntax:

operation ::= `transform.structured.multitile_sizes` $target attr-dict

Emits the IR computing the tile sizes s1 and s2 such that:

• there exists a combination of n tiles of size s1 and m tiles of size s2 that covers the entirety of the iteration space dimension of the target structured op;
• s1, s2 is less than or equal to target_size;
• s1 and s2 are divisible by `divisor.

For example, for a dimension of size 54 with target size 12 and divisor 2, this can emit the IR computing the tile size 10, used for 3 tiles, and 12, used for 2 tiles, totally 103 + 122 = 54. Note that when the divisor does not divide the original dimension size, it is impossible to compute such tile sizes. An assertion is emitted to guard against this in the dynamic case.

Expects the target size and the divisor to be strictly positive. Folds the IR as much as possible, normally obtaining constant sizes and numbers of tiles for a statically known dimension.

This does not consume the target handle and produces three handles each pointing to single-result index-typed operations (which may be arithmetic constant operations) defining the two respective tile sizes and the product of the first tile size with the number of tiles of that size (useful for splitting the iteration space).

This operation composes with the regular tiling when applied per-dimension:

%sz1, %sz2, %split = structured.multitile_sizes %target
                     { target_size = 10, dimension = 1 }
%low, %high = structured.split %target after %split { dimension = 1 }
%tiled_low = structured.tile %low [0, %sz1]
%tiled_high = structured.tile %high [0, %sz2]
%common = merge_handles %tiled_low, %tiled_high

%sz3, %sz4, %split = structured.multitile_size %target
                     { target_size = 42, dimension = 0 }
%sz3r, %sz4r, %splitr = replicate num(%common) %sz3, %sz4, %splitr
structured.split %common after %splitr { dimension = 0 }
// ...

Traits: TransformEachOpTrait

Interfaces: MemoryEffectOpInterface, TransformOpInterface

Attributes: ¶

Operands: ¶

Results: ¶

transform.structured.pad (::mlir::transform::PadOp) ¶

Syntax:

operation ::= `transform.structured.pad` $target attr-dict

Pads the operations pointed to by the target handle using the options provides as operation attributes.

Return modes ¶

This operation ignores non-Linalg ops and drops them in the return. This operation may produce a definiteFailure if the padding fails for any reason. If all the operations referred to by the target PDLOperation pad properly, the transform succeeds. Otherwise the transform silently fails. The return handle points to only the subset of successfully produced padded operations, which can be empty.

Traits: FunctionalStyleTransformOpTrait, TransformEachOpTrait

Interfaces: MemoryEffectsOpInterface, TransformOpInterface

Attributes: ¶

Operands: ¶

Results: ¶

transform.structured.promote (::mlir::transform::PromoteOp) ¶

Syntax:

operation ::= `transform.structured.promote` $target attr-dict

Promotes the specified operands of the target into a separate memory buffer.

At this point, this transform does not allow customizing alloc/dealloc functions nor the behavior on copy in/out operations.

Return modes ¶

This operation applies to a single Linalg op that satisfies the promoteSubviewsPrecondition, otherwise it fails.

If the operations referred to by the target PDLOperation promote properly, the transform succeeds.

When successful, the return handle points to the $target operation that was modified inplace.

Traits: FunctionalStyleTransformOpTrait, TransformEachOpTrait

Interfaces: MemoryEffectsOpInterface, TransformOpInterface

Attributes: ¶

Operands: ¶

Results: ¶

transform.structured.replace (::mlir::transform::ReplaceOp) ¶

Syntax:

operation ::= `transform.structured.replace` $target attr-dict-with-keyword regions

Replace all target payload ops with the single op that is contained in this op’s region. All targets must have zero arguments and must be isolated from above.

This op is for debugging/experiments only.

Return modes ¶

This operation consumes the target handle.

Traits: HasOnlyGraphRegion, IsolatedFromAbove, NoTerminator, SingleBlock

Interfaces: MemoryEffectOpInterface, RegionKindInterface, TransformOpInterface

Operands: ¶

Results: ¶

transform.structured.scalarize (::mlir::transform::ScalarizeOp) ¶

Syntax:

operation ::= `transform.structured.scalarize` $target attr-dict

Indicates that ops of a specific kind in the given function should be scalarized (i.e. their dynamic dimensions tiled by 1).

Return modes: ¶

This operation ignores non-Linalg ops and drops them in the return. This operation produces definiteFailure if the scalarization fails for any reason. If all the operations referred to by the target PDLOperation scalarize properly, the transform succeeds. Otherwise the transform silently fails.

The return handle points to only the subset of successfully produced tiled-by-1 operations, which can be empty.

This operation does not return handles to the tiled loop. We make this design choice because it is hard to know ahead of time the number of loops that will be produced (it depends on the number of dynamic dimensions after multiple transformations have been applied). Loops can always be recovered by navigating from the tiled operations if needed.

Traits: FunctionalStyleTransformOpTrait, TransformEachOpTrait

Interfaces: MemoryEffectsOpInterface, TransformOpInterface

Operands: ¶

Results: ¶

transform.structured.split (::mlir::transform::SplitOp) ¶

Indicates that the given target op should be split into two complementary parts, which combined cover the entire iteration domain of the original op. The split is performed along the iteration space dimension provided as attribute. In case of dimension overflow, the transformation fails. The split is performed at the dimension iterator value specified as either the static split point attribute when it is known at transform IR construction time or as the handle to an operation producing a single index-typed value when it is computed by payload IR. In the latter case, the static split point must be set to ShapedType::kDynamic and the dynamic size handle must point to as many value-producing operations as there are structured operations pointed to by the target handle.

The operation consumes the target handle, but preserves the split point handle if provided. It produces two new handles pointing to the two parts of the structured op after splitting, in the same order as the target operand, with the first handle corresponding to the part with lower iteration space indices.

Interfaces: MemoryEffectOpInterface, TransformOpInterface

Attributes: ¶

Operands: ¶

Results: ¶

transform.structured.split_reduction (::mlir::transform::SplitReductionOp) ¶

Syntax:

operation ::= `transform.structured.split_reduction` $target attr-dict

Indicates that the given target op should be transformed with the splitReduction transformation and split factor provided as attribute.

The splitReduction transformation splits the first single linalg op reduction into a parallel and reduction dimension. A new linalg.generic op is created to perform the rest of the reduction.

The transformation supports different configurations attributes:

• split_factor: the factor by which to split (i.e. the size of the remaining reduction after splitting).
• insert_split_dimension: the dimension in the temporary tensor into which the new parallel dimension is inserted.
• inner_parallel: specifies whether the parallel dimension is before or after the reduction dimension in the splitting op.
• use_scaling_algorithm: whether to use a scaling based formulation that does not create an ExpandShapeOp (default: do not use scaling)
• use_alloc: whether to use an alloc op to allocate the temporary tensor (default: do not use alloc op)

Return modes ¶

This operation ignores non-Linalg ops and drops them in the return. This operation produces definiteFailure if the splitting fails for any reason.

If all the operations referred to by the target PDLOperation split properly, the transform succeeds. Otherwise the transform silently fails. The 4 returned handles points to only the subset of successfully produced computational operations, which can all be empty. This 4 returned handles point to:

• the init op (or tensor_alloc op if use_alloc = true),
• the fill op used to initialize the neutral element,
• the split op and
• the result-combining op.

Example (default: use_scaling_algorithm = false, use_alloc = false): ¶

  %r = linalg.generic {indexing_maps = [affine_map<(d0) -> (d0)>,
                                        affine_map<(d0) -> ()>],
        iterator_types = ["reduction"]}
  ins(%in : tensor<32xf32>)
  outs(%out : tensor<f32>) {
  ^bb0(%arg1: f32, %arg2: f32):
    %y = arith.addf %arg1, %arg2 : f32
    linalg.yield %y : f32
  } -> tensor<f32>

is split into:

  %cst = arith.constant 0.000000e+00 : f32
  %0 = tensor.expand_shape %in [[0, 1]] : tensor<32xf32> into tensor<4x8xf32>
  %1 = tensor.empty() : tensor<4xf32>
  %2 = linalg.fill ins(%cst : f32) outs(%1 : tensor<4xf32>) -> tensor<4xf32>
  %3 = linalg.generic {indexing_maps = [affine_map<(d0, d1) -> (d0, d1)>,
                                        affine_map<(d0, d1) -> (d0)>],
    iterator_types = ["parallel", "reduction"]}
    ins(%0 : tensor<4x8xf32>) outs(%2 : tensor<4xf32>) {
    ^bb0(%arg3: f32, %arg5: f32):
    %5 = arith.addf %arg3, %arg4 : f32
    linalg.yield %5 : f32
  } -> tensor<4xf32>
  %r = linalg.generic {indexing_maps = [affine_map<(d0) -> (d0)>,
                                        affine_map<(d0) -> ()>],
    iterator_types = ["reduction"]}
    ins(%3 : tensor<4xf32>) outs(%out : tensor<f32>) {
    ^bb0(%arg3: f32, %arg4: f32):
    %5 = arith.addf %arg3, %arg4 : f32
    linalg.yield %5 : f32
  } -> tensor<f32>

Example (use_scaling_algorithm = true, use_alloc = true): ¶

Instead of introducing an ExpandShapeOp, this scaling-based implementation rewrites a reduction dimension k into k * split_factor + kk. The dimension kk is added as an extra parallel dimension to the intermediate output tensor at position insert_split_dimension.

Consider a minimal example where k is reduced: O(i, j) += I(i, j, k) Assume i=3, j=5, k=128, split_factor=16 and insert_split_dimension=0. The compute is rewritten as: a. O_i(kk, i, j) += I(i, j, 16 * k + kk) b. O(i, j) += O_i(kk, i, j) The intermediate tensor O_i is of shape (128/16)x3x5 == 8x3x5.

Example: ¶

 %0 = linalg.matmul ins(%A, %B: tensor<16x256xf32>, tensor<256x32xf32>)
   outs(%C: tensor<16x32xf32>) -> tensor<16x32xf32>

Is transformed to:

 #map0 = affine_map<(d0, d1, d2, d3) -> (d0, d2 * 4 + d3)>
 #map1 = affine_map<(d0, d1, d2, d3) -> (d2 * 4 + d3, d1)>
 #map2 = affine_map<(d0, d1, d2, d3) -> (d2, d3)>
 #map3 = affine_map<(d0, d1, d2, d3) -> (d0, d1, d2)>
 #map4 = affine_map<(d0, d1, d2) -> (d0, d1, d2)>
 #map5 = affine_map<(d0, d1, d2) -> (d0, d1)>
 %0 = tensor.empty() : tensor<16x32x64xf32>
 %cst = arith.constant 0.000000e+00 : f32
 %1 = linalg.fill ins(%cst : f32) outs(%0 : tensor<16x32x64xf32>) ->
    tensor<16x32x64xf32>
 %2 = tensor.empty() : tensor<64x4xi1>

 %3 = linalg.generic {indexing_maps = [#map0, #map1, #map2, #map3],
   iterator_types = ["parallel", "parallel", "parallel", "reduction"]}
   ins(%A, %B, %2 : tensor<16x256xf32>, tensor<256x32xf32>, tensor<64x4xi1>)
   outs(%1 : tensor<16x32x64xf32>) {
     ^bb0(%arg3: f32, %arg4: f32, %arg5: i1, %arg6: f32):
       %5 = arith.mulf %arg3, %arg4 : f32
       %6 = arith.addf %arg6, %5 : f32
       linalg.yield %6 : f32
 } -> tensor<16x32x64xf32>

 %4 = linalg.generic {indexing_maps = [#map4, #map5],
   iterator_types = ["parallel", "parallel", "reduction"]}
   ins(%3 : tensor<16x32x64xf32>)
   outs(%C : tensor<16x32xf32>) {
     ^bb0(%arg3: f32, %arg4: f32):
       %5 = arith.addf %arg3, %arg4 : f32
       linalg.yield %5 : f32
 } -> tensor<16x32xf32>

 return %4 : tensor<16x32xf32>

Traits: FunctionalStyleTransformOpTrait, TransformEachOpTrait

Interfaces: MemoryEffectsOpInterface, TransformOpInterface

Attributes: ¶

Operands: ¶

Results: ¶

transform.structured.tile (::mlir::transform::TileOp) ¶

Indicates that the given target op should be tiled with the given sizes. This transform generates a loop nest with a smaller (“tiled”) target operation in its body. Currently limited to LinalgOps.

Tile sizes may be known at transformation time, in which case they are expected to be provided in the static_size attribute, or not, in which case the tile value must be computed by the payload IR and the handle to the operation computing it must be provided through dynamic_sizes. When the sizes are not known statically, the corresponding entry in the static_sizes attribute must be set to ShapedType::kDynamic. Only the dynamic sizes must be provided in dynamic_sizes, i.e., there should be as many handles as ShapedType::kDynamic values in the static_sizes attribute. A static size of 0 indicates that the dimension should not be tiled. No loop will be generated for such dimensions. If all tile sizes are 0, this transform is effectively a no-op.

This op returns handles to the tiled op (in the generated loop nest) and the generated loops. The number of loops is the number of tile sizes that are statically known to be non-zero.

Return modes ¶

On success, the resulting handles are associated with co-indexed lists of tiled operations and loops around them.

This operation only supports Linalg ops and produces a silenceable failure if the input contains any non-Linalg ops. The ops preceding it in the list associated with the target handle will have been tiled.

This operation produces a silenceable failure if the dynamic_sizes handles are associated with lists of payload operations of a size different than that of the list associated with the target handle.

If the internal implementation of tiling for any of the operations fails, produces a definite failure.

Interfaces: MemoryEffectOpInterface, TransformOpInterface

Attributes: ¶

Operands: ¶

Results: ¶

transform.structured.tile_reduction_using_foreach_thread (::mlir::transform::TileReductionUsingForeachThreadOp) ¶

Syntax:

operation ::= `transform.structured.tile_reduction_using_foreach_thread` $target attr-dict

Tile a PartialReductionOpInterface op to a tiled scf.foreach_thread doing partial reduction.

This transformation tiles the target along the reduction dimensions. It creates a tensor initialized with the identity value. Then it creates a scf.foreach_thread loops with the number threads given by num_threads. The op is tiled op with a size equal to floordiv(size, num_threads). All the partial reduction value is are parallel inserted to create a new tensor. After the loop a merge operation is created to do a final reduction with the partial reductions tensor. If an extra tile_sizes parameter is passed the tiles are cyclically distributed on the threads of the scf.foreach_threads loop.

Return modes ¶

This 3 returned handles point to:

• the fill op used to initialize the neutral element,
• the parallel tiled op and
• the result-combining op.

Example: ¶

  %red = linalg.generic {indexing_maps = [affine_map<(d0, d1) -> (d0, d1)>,
                                          affine_map<(d0, d1) -> (d0)>],
  iterator_types = ["parallel", "reduction"]}
  ins(%arg0 : tensor<?x?xf32>)
  outs(%out : tensor<?xf32>) {
    ^bb0(%arg7: f32, %arg9: f32):
    %1 = arith.addf %arg7, %arg9 : f32
    linalg.yield %1 : f32
  } -> tensor<?xf32>
  return %red : tensor<?xf32>

is transformed into:

  %0 = tensor.empty(%dim_1) : tensor<?x5xf32>
  %1 = linalg.fill ins(%cst : f32) outs(%0 : tensor<?x5xf32>) -> tensor<?x5xf32>
  %2 = scf.foreach_thread (%arg2) in (%c5) shared_outs(%arg3 = %1) -> (tensor<?x5xf32>) {
    %4 = affine.min #map(%arg2)[%dim_0]
    %5 = affine.max #map1(%4)
    %extracted_slice = tensor.extract_slice %arg3[0, %arg2] [%dim, 1] [1, 1] : tensor<?x5xf32> to tensor<?xf32>
    %6 = affine.apply #map2(%arg2)[%dim_0]
    %extracted_slice_2 = tensor.extract_slice %arg0[0, %6] [%dim, %5] [1, 1] : tensor<?x?xf32> to tensor<?x?xf32>
    %extracted_slice_3 = tensor.extract_slice %extracted_slice[0] [%dim] [1] : tensor<?xf32> to tensor<?xf32>
    %7 = linalg.generic {indexing_maps = [#map3, #map4], iterator_types = ["parallel", "reduction"]} ins(%extracted_slice_2 : tensor<?x?xf32>) outs(%extracted_slice_3 : tensor<?xf32>) {
    ^bb0(%in: f32, %out: f32):
      %9 = arith.addf %in, %out : f32
      linalg.yield %9 : f32
    } -> tensor<?xf32>
    scf.foreach_thread.perform_concurrently {
      tensor.parallel_insert_slice %7 into %arg3[0, %arg2] [%dim, 1] [1, 1] : tensor<?xf32> into tensor<?x5xf32>
    }
  } {thread_dim_mapping = []}
  %3 = linalg.generic {indexing_maps = [#map3, #map4], iterator_types = ["parallel", "reduction"]} ins(%2 : tensor<?x5xf32>) outs(%arg1 : tensor<?xf32>) {
  ^bb0(%in: f32, %out: f32):
    %4 = arith.addf %in, %out : f32
    linalg.yield %4 : f32
  } -> tensor<?xf32>

Traits: FunctionalStyleTransformOpTrait, TransformEachOpTrait

Interfaces: MemoryEffectsOpInterface, TransformOpInterface

Attributes: ¶

Operands: ¶

Results: ¶

transform.structured.tile_reduction_using_scf (::mlir::transform::TileReductionUsingScfOp) ¶

Syntax:

operation ::= `transform.structured.tile_reduction_using_scf` $target attr-dict

Indicates that the given target op should be transformed with the tileReduction transformation with the tile size provided as attribute.

This transformation tiles the target along the reduction dimensions. It creates a tensor initialized with the identity value. Then it creates nested loops with a parallel version of target op inside. The parallel op dimensions are less or equal to the tile size passed by user. After the loop a merge operation is created to do a final reduction with the partial reductions. The initial tensor always uses the tile size dimension. This may overallocate if the tile size is greater than the reduction dimension.

Return modes ¶

This 3 returned handles point to:

• the fill op used to initialize the neutral element,
• the parallel tiled op and
• the result-combining op.

Example: ¶

  %red = linalg.generic {indexing_maps = [affine_map<(d0, d1) -> (d0, d1)>,
                                          affine_map<(d0, d1) -> (d0)>],
  iterator_types = ["parallel", "reduction"]}
  ins(%arg0 : tensor<?x?xf32>)
  outs(%out : tensor<?xf32>) {
    ^bb0(%arg7: f32, %arg9: f32):
    %1 = arith.addf %arg7, %arg9 : f32
    linalg.yield %1 : f32
  } -> tensor<?xf32>
  return %red : tensor<?xf32>

is transformed into:

  %0 = tensor.empty(%dim_1) : tensor<?x5xf32>
  %1 = linalg.fill ins(%cst : f32) outs(%0 : tensor<?x5xf32>) -> tensor<?x5xf32>
  %2 = scf.for %arg2 = %c0 to %dim_0 step %c5 iter_args(%arg3 = %1) -> (tensor<?x5xf32>) {
    %extracted_slice = tensor.extract_slice %1[0, 0] [%dim, 5] [1, 1] : tensor<?x5xf32> to tensor<?x5xf32>
    %extracted_slice_2 = tensor.extract_slice %arg0[0, %arg2] [%dim, 5] [1, 1] : tensor<?x?xf32> to tensor<?x5xf32>
    %4 = linalg.generic {indexing_maps = [affine_map<(d0, d1) -> (d0, d1)>,
                                          affine_map<(d0, d1) -> (d0, d1)>],
    iterator_types = ["parallel", "parallel"]}
    ins(%extracted_slice_2 : tensor<?x5xf32>)
    outs(%extracted_slice : tensor<?x5xf32>) {
    ^bb0(%in: f32, %out: f32):
      %5 = arith.addf %in, %out : f32
      linalg.yield %5 : f32
    } -> tensor<?x5xf32>
    %dim_3 = tensor.dim %1, %c0 : tensor<?x5xf32>
    %inserted_slice = tensor.insert_slice %4 into %arg3[0, 0] [%dim_3, 5] [1, 1] : tensor<?x5xf32> into tensor<?x5xf32>
    scf.yield %inserted_slice : tensor<?x5xf32>
  }
  %3 = linalg.generic {indexing_maps = [affine_map<(d0, d1) -> (d0, d1)>,
                                        affine_map<(d0, d1) -> (d0)>],
  iterator_types = ["parallel", "reduction"]}
  ins(%2 : tensor<?x5xf32>)
  outs(%arg1 : tensor<?xf32>) {
  ^bb0(%in: f32, %out: f32):
    %4 = arith.addf %in, %out : f32
    linalg.yield %4 : f32
  } -> tensor<?xf32>

Traits: FunctionalStyleTransformOpTrait, TransformEachOpTrait

Interfaces: MemoryEffectsOpInterface, TransformOpInterface

Attributes: ¶

Operands: ¶

Results: ¶

transform.structured.tile_to_foreach_thread_op (::mlir::transform::TileToForeachThreadOp) ¶

Syntax:

operation ::= `transform.structured.tile_to_foreach_thread_op` $target oilist(
              `num_threads` custom<DynamicIndexList>($num_threads,
              $static_num_threads) |
              `tile_sizes` custom<DynamicIndexList>($tile_sizes,
              $static_tile_sizes))
              (`(` `mapping` `=` $mapping^ `)`)? attr-dict

Tile a TilingInterface op to a tiled scf.foreach_thread.

Tiling is applied by either specifying num_threads or tile_size. If num_threads is specified, then the tile size for each dimension i is calculated dynamically via ceilDiv(dimSize[i], num_threads[i]). num_threads and tile_size can be either static index attributes or SSA values of PDL operation handle type (or a mix thereof). Operation handles must be mapped to exactly one op that has exactly one result of index type.

Static zero tile sizes indicate that the dimension is not tiled and can be thought of as tiling by the full size of data.

It is the user’s responsibility to ensure that num_threads/tile_sizes is a valid tiling specification (i.e. that only tiles parallel dimensions, e.g. in the Linalg case).

If non-empty, the mapping is added as an attribute to the resulting scf.foreach_thread.

Note: tile_sizes and num_threads are variadic. Each tile size/number of threads can be an index attribute or a transform handle that is mapped to exactly one payload op with exactly one index result.

Return modes ¶

This operation ignores ops that do not implement the TilingInterface and drops them in the return.

If all the operations referred to by the target PDLOperation tile successfully, the transform succeeds. Otherwise the transform silently fails.

The two returned handles point to only the subset of successfully produced tiled operations, which can all be empty.

These two returned handles point to:

• the new scf.foreach_thread op,
• the tiled op that implements TilingInterface.

Example using num_threads ¶

%0 = pdl_match @match_matmul in %arg1
%3:2 = transform.structured.tile_to_foreach_thread_op %0 num_threads [10, 20]

Example using tile_sizes ¶

%0 = pdl_match @match_matmul in %arg1
%sz = pdl_match @match_size_op in %arg1
%3:2 = transform.structured.tile_to_foreach_thread_op %0 tile_sizes [0, %sz, 20]

Traits: AttrSizedOperandSegments

Interfaces: MemoryEffectOpInterface, TransformOpInterface

Attributes: ¶

Operands: ¶

Results: ¶

transform.structured.tile_to_scf_for (::mlir::transform::TileToScfForOp) ¶

Indicates that the given target op should be tiled with the given sizes. This transform generates a loop nest with a smaller (“tiled”) target operation in its body. The target must implement TilingInterface.

Tile sizes may be known at transformation time, in which case they are expected to be provided in the static_size attribute, or not, in which case the tile value must be computed by the payload IR and the handle to the operation computing it must be provided through dynamic_sizes. When the sizes are not known statically, the corresponding entry in the static_sizes attribute must be set to ShapedType::kDynamic. Only the dynamic sizes must be provided in dynamic_sizes, i.e., there should be as many handles as ShapedType::kDynamic values in the static_sizes attribute. A static size of 0 indicates that the dimension should not be tiled. No loop will be generated for such dimensions. If all tile sizes are 0, this transform is effectively a no-op.

This op returns handles to the tiled op (in the generated loop nest) and the generated loops. The number of loops is the number of tile sizes that are statically known to be non-zero.

Return modes ¶

On success, the resulting handles are associated with co-indexed lists of tiled operations and loops around them.

This operation only supports TilingInterface ops and produces a silenceable failure if the input contains any non-TilingInterface ops. The ops preceding it in the list associated with the target handle will have been tiled.

This operation produces a silenceable failure if the dynamic_sizes handles are associated with lists of payload operations of a size different than that of the list associated with the target handle.

If the internal implementation of tiling for any of the operations fails, produces a definite failure.

Interfaces: MemoryEffectOpInterface, TransformOpInterface

Attributes: ¶

Operands: ¶

Results: ¶

transform.structured.vectorize (::mlir::transform::VectorizeOp) ¶

Syntax:

operation ::= `transform.structured.vectorize` $target attr-dict

Indicates that the given target op all the ops it contains should be vectorized with the configuration specified by the attributes of this op. This vectorization only handles structured ops that operate on shaped types and does not vectorize loops or straight-line. Internally, it applies a set of rewrite patterns, some of which enable vectorization and some of which clean up the results. Therefore, it can only be applied to an op with the “isolated from above property”. If finer granularity is required, it can be achieved by outlining the target part of the payload IR into, e.g., a function, performing the transformation, and inlining it back. This transformation only fails if the entire pattern rewriting failed, i.e., it does not fail when no ops were vectorized.

Note that this transformation is invalidating the handles to any payload IR operation that is contained inside the vectorization target.

This transformation supports the following attributes:

• vectorize_padding: a UnitAttr to activate the vectorization of tensor.pad ops. Different pipelines may prefer to lower such ops to loops.
• disable_multi_reduction_to_contract_patterns: a UnitAttr to deactivate the rewrite of vector.multi_reduction to vector.contract. This is intended to be used in tests only.
• disable_transfer_permutation_map_lowering_patterns: a UnitAttr to deactivate the rewrite of vector.transfer with permutation maps into explicit vector.transpose operations. This is intended to be used in tests only but may be promotoed to a first class attribute in the future.

Return modes: ¶

This operation produces definiteFailure if vectorization fails for any reason. The operation always returns the handle to the target op that is expected to be isolated from above.

Traits: FunctionalStyleTransformOpTrait, TransformEachOpTrait

Interfaces: MemoryEffectsOpInterface, TransformOpInterface

Attributes: ¶

Operands: ¶

Results: ¶

TransformTypeInterface (TransformTypeInterface) ¶

Types that can be used for Transform dialect handle values. Such types define the properties of Payload IR operations associated with the handle. A user of such a handle can assume that these properties have been verified for any Payload IR operation associated with it.

Methods: ¶

checkPayload ¶

::mlir::DiagnosedSilenceableFailure checkPayload(::mlir::Location loc, ::mlir::ArrayRef<::mlir::Operation *> payload);

Checks if the given list of associated Payload IR operations satisfy the conditions defined by this type. If not, produces a silenceable error at the specified location.

NOTE: This method must be implemented by the user.

TransformOpInterface (TransformOpInterface) ¶

This interface is to be implemented by operations that identify transformations to be performed on other operations. The former are referred to as transform IR operations. The latter are referred to as payload IR operations. Such transform IR operations provide a fine-grain control mechanism over how transformations are applied by using and defining transform IR values, referred to as handles, that correspond to sets of operations in the payload IR. Transformations are applied starting from the operations identified by handles, but may affect other operations as well. Further restrictions may be imposed by flows that rely on transform IR operations to control transformations.

Methods: ¶

apply ¶

::mlir::DiagnosedSilenceableFailure apply(::mlir::transform::TransformResults &transformResults, ::mlir::transform::TransformState &state);

Applies the transformation represented by the current operation. This accepts as arguments the object that must be populated with results of the current transformation and a transformation state object that can be used for queries, e.g., to obtain the list of operations on which the transformation represented by the current op is targeted. Returns a special status object indicating whether the transformation succeeded or failed, and, if it failed, whether the failure is recoverable or not.

NOTE: This method must be implemented by the user.