'iree_input' Dialect

Public ops/type/attributes legal for input to IREE's compiler.

IREE's compiler allows as input a number of common dialects. This dialect contains structural and unique ops that do not exist elsewhere or that IREE has an interest in maintaining as a stable set.

The contents of this dialect often mirror various constructs in IREE's internal implementation. The focus here is on simplicity and stability over time. Generally, this dialect does not use "advanced" features and should be broadly source compatible over a range of LLVM versions. There are of course, limits, and source-compatibility is not guaranteed, since LLVM/MLIR's API surface is itself unstable.

• 'iree_input' Dialect
  • Operations
    • Buffer and buffer view ops
      • iree_input.buffer.subspan (Input::BufferSubspanOp)
      • iree_input.buffer_view.create (Input::BufferViewCreateOp)
      • iree_input.buffer_view.dim (Input::BufferViewDimOp)
      • iree_input.buffer_view.rank (Input::BufferViewRankOp)
    • Byte buffer ops
      • iree_input.byte_buffer.constant (Input::ByteBufferConstantOp)
    • Compiler hint ops
      • iree_input.optimization_barrier (Input::OptimizationBarrierOp)
    • Dispatch ops
      • iree_input.dispatch (Input::DispatchOp)
    • Executable source ops
      • iree_input.executable.export (Input::ExecutableExportOp)
      • iree_input.executable.source_end (Input::ExecutableSourceEndOp)
      • iree_input.executable.source (Input::ExecutableSourceOp)
    • Global variable ops
      • iree_input.global.address (Input::GlobalAddressOp)
      • iree_input.global.load.indirect (Input::GlobalLoadIndirectOp)
      • iree_input.global.load (Input::GlobalLoadOp)
      • iree_input.global (Input::GlobalOp)
      • iree_input.global.store.indirect (Input::GlobalStoreIndirectOp)
      • iree_input.global.store (Input::GlobalStoreOp)
    • Mutable list ops
      • iree_input.list.create (Input::ListCreateOp)
      • iree_input.list.get (Input::ListGetOp)
      • iree_input.list.resize (Input::ListResizeOp)
      • iree_input.list.set (Input::ListSetOp)
      • iree_input.list.size (Input::ListSizeOp)
    • Pseudo ops for conversion support
      • iree_input.tensor.export (Input::TensorExportOp)
      • iree_input.tensor.import (Input::TensorImportOp)
    • Tensor ops
      • iree_input.tensor.bitcast (Input::TensorBitCastOp)
      • iree_input.tensor.clone (Input::TensorCloneOp)
      • iree_input.tensor.load (Input::TensorLoadOp)
      • iree_input.tensor.reshape (Input::TensorReshapeOp)
      • iree_input.tensor.slice (Input::TensorSliceOp)
      • iree_input.tensor.splat (Input::TensorSplatOp)
      • iree_input.tensor.store (Input::TensorStoreOp)
      • iree_input.tensor.trace (Input::TensorTraceOp)
      • iree_input.tensor.update (Input::TensorUpdateOp)
    • Utility ops
      • iree_input.align (Input::AlignOp)
      • iree_input.null (Input::NullOp)
    • Workgroup dispatch ops
      • iree_input.dispatch.workgroup.count (Input::DispatchWorkgroupCountOp)
      • iree_input.dispatch.workgroup.id (Input::DispatchWorkgroupIDOp)
      • iree_input.dispatch.workgroup.size (Input::DispatchWorkgroupSizeOp)
  • Attributes
    • DescriptorSetBindingAttr
    • DescriptorSetLayoutAttr
    • DescriptorTypeAttr
    • DeviceTargetAttr
    • ExecutableObjectAttr
    • ExecutableObjectsAttr
    • ExecutableTargetAttr
    • PipelineLayoutAttr
  • Type constraints
    • list
  • Types
    • BufferType
    • BufferViewType
    • ByteBufferType
    • ListType
    • PtrType
    • VariantType

Operations

Buffer and buffer view ops

iree_input.buffer.subspan (Input::BufferSubspanOp)

Buffer subspan operation

Syntax:

operation ::= `iree_input.buffer.subspan` `<` $source_buffer `:` type($source_buffer) `>`
              `` `[` $source_offset `,` $length `]`
              `:` type($result)
              attr-dict-with-keyword

Returns a reference to a subspan of the buffer.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface), OpAsmOpInterface

Effects: MemoryEffects::Effect{}

Operands:

Operand	Description
source_buffer	Buffer is an untyped bag of bits with no shape or dtype
source_offset	index
length	index

Results:

Result	Description
result	Buffer is an untyped bag of bits with no shape or dtype

iree_input.buffer_view.create (Input::BufferViewCreateOp)

Buffer view reference initializer

Syntax:

operation ::= `iree_input.buffer_view.create` `buffer` `(` $source_buffer `:` type($source_buffer) `)`
              `` `[` $source_offset `,` $source_length `]`
              `shape` `(` `[` $shape `]` `)`
              `type` `(` $element_type `)`
              `encoding` `(` $encoding_type `)`
              `:` type($result)
              attr-dict-with-keyword

Creates a reference to a buffer with a particular shape and element type. The buffer is not copied and both the original and view references must be synchronized. This makes it easier to associate commonly-carried metadata along with the contents.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Operand	Description
source_buffer	Buffer is an untyped bag of bits with no shape or dtype
source_offset	index
source_length	index
element_type	32-bit signless integer
encoding_type	32-bit signless integer
shape	variadic of index

Results:

Result	Description
result	View into a buffer, with runtime shape and element type

iree_input.buffer_view.dim (Input::BufferViewDimOp)

Buffer view dimension value query

Syntax:

operation ::= `iree_input.buffer_view.dim` $buffer_view `,` $index attr-dict `:` type($result)

Returns the value of the given dimension.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Attribute	MLIR Type	Description
index	::mlir::IntegerAttr	index attribute

Operands:

Operand	Description
buffer_view	View into a buffer, with runtime shape and element type

Results:

Result	Description
result	index

iree_input.buffer_view.rank (Input::BufferViewRankOp)

Buffer view rank query

Syntax:

operation ::= `iree_input.buffer_view.rank` $buffer_view attr-dict `:` type($result)

Returns the rank of the buffer view.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Operand	Description
buffer_view	View into a buffer, with runtime shape and element type

Results:

Result	Description
result	index

Byte buffer ops

iree_input.byte_buffer.constant (Input::ByteBufferConstantOp)

Constant host-side byte buffer

Syntax:

operation ::= `iree_input.byte_buffer.constant` ($name^)? attr-dict `:` type($result) `=` $value

Defines a compile-time byte buffer based on the given attribute value. The attribute will be serialized into the canonical IREE format for the chosen host target.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface), OpAsmOpInterface

Effects: MemoryEffects::Effect{}

Attributes:

Attribute	MLIR Type	Description
name	::mlir::StringAttr	string attribute
value	::mlir::StringAttr	string attribute
alignment	::mlir::IntegerAttr	index attribute
mime_type	::mlir::StringAttr	string attribute

Results:

Result	Description
result	a reference counted byte buffer

Compiler hint ops

iree_input.optimization_barrier (Input::OptimizationBarrierOp)

Prevents compiler optimizations across a value.

Syntax:

operation ::= `iree_input.optimization_barrier` attr-dict
              ($operands^ `:` type($operands))?

Wraps any operands in an unoptimizable identity to prevent its results from being folded. It will be dropped during the final step in compilation and has no effect at runtime.

Traits: SameOperandsAndResultType

Operands:

Operand	Description
operands	variadic of any type

Results:

Result	Description
results	variadic of any type

Dispatch ops

iree_input.dispatch (Input::DispatchOp)

A dispatch of an executable across a grid

Syntax:

operation ::= `iree_input.dispatch` $entry_point
              (`[` $workload^ `]`)? ``
              `(` $arguments `)` attr-dict `:`
              custom<ShapedFunctionType>(ref($arguments),
              type($arguments), $argument_dims,
              type($results), $result_dims,
              $tied_operands)

Traits: AlwaysSpeculatableImplTrait, AttrSizedOperandSegments

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface), SymbolUserOpInterface, TiedOpInterface

Effects: MemoryEffects::Effect{}

Attributes:

Attribute	MLIR Type	Description
entry_point	::mlir::SymbolRefAttr	symbol reference attribute
tied_operands	::mlir::ArrayAttr	64-bit integer array attribute

Operands:

Operand	Description
workload	variadic of index
arguments	variadic of any type
argument_dims	variadic of index
result_dims	variadic of index

Results:

Result	Description
results	variadic of any type

Executable source ops

iree_input.executable.export (Input::ExecutableExportOp)

Executable entry point declaration

Syntax:

operation ::= `iree_input.executable.export` custom<SymbolVisibility>($sym_visibility)
              $sym_name
              `ordinal` `(` $ordinal `)`
              `layout` `(` $layout `)`
              attr-dict-with-keyword

Traits: HasParent<IREE::Input::ExecutableSourceOp>, IsolatedFromAbove

Interfaces: Symbol

Attributes:

Attribute	MLIR Type	Description
sym_visibility	::mlir::StringAttr	string attribute
sym_name	::mlir::StringAttr	string attribute
ordinal	::mlir::IntegerAttr	size_t
layout	::mlir::iree_compiler::IREE::Input::PipelineLayoutAttr	executable entry point layout specification
workgroup_size	::mlir::ArrayAttr	index array attribute
subgroup_size	::mlir::IntegerAttr	size_t
workgroup_local_memory	::mlir::IntegerAttr	index attribute

iree_input.executable.source_end (Input::ExecutableSourceEndOp)

Terminator pseudo-op for the executable source op

Syntax:

operation ::= `iree_input.executable.source_end` attr-dict

Traits: HasParent<IREE::Input::ExecutableSourceOp>, Terminator

iree_input.executable.source (Input::ExecutableSourceOp)

Generic source contents of an executable op

Syntax:

operation ::= `iree_input.executable.source` custom<SymbolVisibility>($sym_visibility)
              $sym_name
              attr-dict-with-keyword
              ``
              regions

Traits: IsolatedFromAbove, SingleBlockImplicitTerminator<IREE::Input::ExecutableSourceEndOp>, SingleBlock, SymbolTable

Interfaces: Symbol

Attributes:

Attribute	MLIR Type	Description
sym_visibility	::mlir::StringAttr	string attribute
sym_name	::mlir::StringAttr	string attribute
objects	::mlir::iree_compiler::IREE::Input::ExecutableObjectsAttr	target-specific object file references

Global variable ops

iree_input.global.address (Input::GlobalAddressOp)

Returns an address reference to a global

Syntax:

operation ::= `iree_input.global.address` $global attr-dict `:` type($result)

Returns the address of a global as a typed reference. Can be used with the global load and store indirect ops.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Attribute	MLIR Type	Description
global	FlatSymbolRefAttr	symbol reference attribute

Results:

Result	Description
result	ranked tensor of any type values or index or signless integer or floating-point or complex-type

iree_input.global.load.indirect (Input::GlobalLoadIndirectOp)

Loads a value from a global variable

Syntax:

operation ::= `iree_input.global.load.indirect` $global attr-dict `:` type($global) `->` type($result)

Returns a copy of the global value.

Operands:

Operand	Description
global	ranked tensor of any type values or index or signless integer or floating-point or complex-type

Results:

Result	Description
result	any type

iree_input.global.load (Input::GlobalLoadOp)

Loads a value from a global variable

Syntax:

operation ::= `iree_input.global.load` $global attr-dict `:` type($result)

Returns a copy of the global value.

Interfaces: SymbolUserOpInterface

Attributes:

Attribute	MLIR Type	Description
global	FlatSymbolRefAttr	symbol reference attribute

Results:

Result	Description
result	any type

iree_input.global (Input::GlobalOp)

Stateful global variable declaration

Syntax:

operation ::= `iree_input.global` custom<SymbolVisibility>($sym_visibility)
              (`mutable` $is_mutable^)?
              $sym_name
              attr-dict
              (`initializer` `(` $initializer^ `)`)?
              custom<TypeOrAttr>($type, $initial_value)

Declares a global variable that maintains its value across invocations. The value is tied to the execution context of the module and different contexts will have different global storage.

Interfaces: Symbol

Attributes:

Attribute	MLIR Type	Description
sym_visibility	::mlir::StringAttr	string attribute
sym_name	::mlir::StringAttr	string attribute
type	::mlir::TypeAttr	any type attribute
is_mutable	::mlir::UnitAttr	unit attribute
initializer	::mlir::FlatSymbolRefAttr	flat symbol reference attribute
initial_value	::mlir::TypedAttr	TypedAttr instance

iree_input.global.store.indirect (Input::GlobalStoreIndirectOp)

Stores a value into a global variable

Syntax:

operation ::= `iree_input.global.store.indirect` $value `,` $global attr-dict `:` type($value) `->` type($global)

Stores a copy of the value into a global.

Operands:

Operand	Description
value	any type
global	ranked tensor of any type values or index or signless integer or floating-point or complex-type

iree_input.global.store (Input::GlobalStoreOp)

Stores a value into a global variable

Syntax:

operation ::= `iree_input.global.store` $value `,` $global attr-dict `:` type($value)

Stores a copy of the value into a global.

Interfaces: SymbolUserOpInterface

Attributes:

Attribute	MLIR Type	Description
global	FlatSymbolRefAttr	symbol reference attribute

Operands:

Operand	Description
value	any type

Mutable list ops

iree_input.list.create (Input::ListCreateOp)

Creates a new empty list

Syntax:

operation ::= `iree_input.list.create` ($initial_capacity^)? attr-dict `:` type($result)

Creates a new empty list with an optional initial capacity.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, MemoryEffectOpInterface (MemoryEffectOpInterface), NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Allocate on ::mlir::SideEffects::DefaultResource}, MemoryEffects::Effect{}

Operands:

Operand	Description
initial_capacity	index

Results:

Result	Description
result	list

iree_input.list.get (Input::ListGetOp)

Element accessor

Syntax:

operation ::= `iree_input.list.get` $list `[` $index `]` attr-dict `:` type($list) `->` type($result)

Returns the value of the element at the given index. Note that the value may be null if the element is null or the type does not match.

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Read on ::mlir::SideEffects::DefaultResource}

Operands:

Operand	Description
list	list
index	index

Results:

Result	Description
result	any type

iree_input.list.resize (Input::ListResizeOp)

Resizes the list to a new count in elements

Syntax:

operation ::= `iree_input.list.resize` operands attr-dict `:` type($list)

Resizes the list to contain new_size elements. This will either truncate the list if the existing size is greater than new_size or extend the list with the default list value of the element type.

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::SideEffects::DefaultResource}

Operands:

Operand	Description
list	list
new_size	index

iree_input.list.set (Input::ListSetOp)

Element mutator

Syntax:

operation ::= `iree_input.list.set` $list `[` $index `]` `,` $value attr-dict `:` type($list) `,` type($value)

Sets the element at the given index to the new value.

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::SideEffects::DefaultResource}

Operands:

Operand	Description
list	list
index	index
value	any type

iree_input.list.size (Input::ListSizeOp)

The size of the list in elements

Syntax:

operation ::= `iree_input.list.size` operands attr-dict `:` type($list)

Returns the current size of the list in elements.

Interfaces: InferTypeOpInterface, MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Read on ::mlir::SideEffects::DefaultResource}

Operands:

Operand	Description
list	list

Results:

Result	Description
result	index

Pseudo ops for conversion support

iree_input.tensor.export (Input::TensorExportOp)

Exports a tensor to a Buffer(View), capturing dynamic dims

Syntax:

operation ::= `iree_input.tensor.export` $source `:` type($source) (`{` $source_dims^ `}`)? `->` type($target)
              attr-dict-with-keyword

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Operand	Description
source	ranked tensor of any type values
source_dims	variadic of index

Results:

Result	Description
target	Buffer is an untyped bag of bits with no shape or dtype or View into a buffer, with runtime shape and element type

iree_input.tensor.import (Input::TensorImportOp)

Imports a Buffer(View) to a tensor, providing dynamic dims

Syntax:

operation ::= `iree_input.tensor.import` $source `:` type($source) `->` type($target) (`{` $target_dims^ `}`)?
              attr-dict-with-keyword

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Operand	Description
source	Buffer is an untyped bag of bits with no shape or dtype or View into a buffer, with runtime shape and element type
target_dims	variadic of index

Results:

Result	Description
target	ranked tensor of any type values

Tensor ops

iree_input.tensor.bitcast (Input::TensorBitCastOp)

Bitcasts a tensor

Syntax:

operation ::= `iree_input.tensor.bitcast` $source `:`
              type($source) (`{` $source_dims^ `}`)? `->`
              type($result) (`{` $result_dims^ `}`)?
              attr-dict-with-keyword

Bitcasts a tensor to a new shape without modifying the contents.

Traits: AlwaysSpeculatableImplTrait, AttrSizedOperandSegments

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Operand	Description
source	ranked tensor of any type values
source_dims	variadic of index
result_dims	variadic of index

Results:

Result	Description
result	ranked tensor of any type values

iree_input.tensor.clone (Input::TensorCloneOp)

Performs a full tensor clone operation

Syntax:

operation ::= `iree_input.tensor.clone` $operand `:` type($result) (`{` $operand_dims^ `}`)?
              attr-dict-with-keyword

Clones the input tensor into an identical output tensor.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Operand	Description
operand	ranked tensor of any type values
operand_dims	variadic of index

Results:

Result	Description
result	ranked tensor of any type values

iree_input.tensor.load (Input::TensorLoadOp)

Loads a value from a tensor element

Syntax:

operation ::= `iree_input.tensor.load` $source (`[` $indices^ `]`)? `:`
              type($source) (`{` $source_dims^ `}`)?
              attr-dict-with-keyword

Returns the element at the given location from within the tensor.

Traits: AlwaysSpeculatableImplTrait, AttrSizedOperandSegments

Interfaces: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Operand	Description
source	ranked tensor of any type values
source_dims	variadic of index
indices	variadic of index

Results:

Result	Description
result	index or signless integer or floating-point or complex-type or vector of any type values

iree_input.tensor.reshape (Input::TensorReshapeOp)

Reshapes a tensor

Syntax:

operation ::= `iree_input.tensor.reshape` $source `:`
              type($source) (`{` $source_dims^ `}`)? `->`
              type($result) (`{` $result_dims^ `}`)?
              attr-dict-with-keyword

Reshapes a tensor to a new shape without modifying the contents.

Traits: AlwaysSpeculatableImplTrait, AttrSizedOperandSegments

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Operand	Description
source	ranked tensor of any type values
source_dims	variadic of index
result_dims	variadic of index

Results:

Result	Description
result	ranked tensor of any type values

iree_input.tensor.slice (Input::TensorSliceOp)

Slices out a subregion of a tensor

Syntax:

operation ::= `iree_input.tensor.slice` $source `[` $start_indices `for` $lengths `]` `:`
              type($source) (`{` $source_dims^ `}`)? `->`
              type($result) (`{` $result_dims^ `}`)?
              attr-dict-with-keyword

Clones a subregion of a tensor.

Traits: AlwaysSpeculatableImplTrait, AttrSizedOperandSegments

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Operand	Description
source	ranked tensor of any type values
source_dims	variadic of index
start_indices	variadic of index
lengths	variadic of index
result_dims	variadic of index

Results:

Result	Description
result	ranked tensor of any type values

iree_input.tensor.splat (Input::TensorSplatOp)

Splats a value into a shaped tensor

Syntax:

operation ::= `iree_input.tensor.splat` $value `:` type($result) (`{` $result_dims^ `}`)?
              attr-dict-with-keyword

Returns a tensor initialized to the given primitive value.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Operand	Description
value	index or signless integer or floating-point or complex-type
result_dims	variadic of index

Results:

Result	Description
result	ranked tensor of any type values

iree_input.tensor.store (Input::TensorStoreOp)

Stores a value into a tensor element

Syntax:

operation ::= `iree_input.tensor.store` $value `,` $target (`[` $indices^ `]`)? `:`
              type($target) (`{` $target_dims^ `}`)?
              attr-dict-with-keyword

Returns a tensor with the element at the given index set to the given value.

Traits: AlwaysSpeculatableImplTrait, AttrSizedOperandSegments

Interfaces: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Operand	Description
value	index or signless integer or floating-point or complex-type or vector of any type values
target	ranked tensor of any type values
target_dims	variadic of index
indices	variadic of index

Results:

Result	Description
result	ranked tensor of any type values

iree_input.tensor.trace (Input::TensorTraceOp)

Traces one or more tensor values at runtime

Syntax:

operation ::= `iree_input.tensor.trace` $key `=` `[`
              custom<ShapedOperandList>($values, type($values), $value_dims)
              `]` attr-dict-with-keyword

Traces out to a runtime trace sink (console, log file, etc) the given tensors. The key is arbitrary and can be used for identifying the set of values being traced.

Traits: AttrSizedOperandSegments

Attributes:

Attribute	MLIR Type	Description
key	::mlir::StringAttr	string attribute

Operands:

Operand	Description
values	variadic of ranked tensor of any type values
value_dims	variadic of index

iree_input.tensor.update (Input::TensorUpdateOp)

Updates a tensor with the contents of another tensor

Syntax:

operation ::= `iree_input.tensor.update` $update `,` $target `[` $start_indices `]` `:`
              type($update) (`{` $update_dims^ `}`)? `->`
              custom<ShapedTiedResult>(type($result), $target_dims)
              attr-dict-with-keyword

Updates the target tensor with the contents of the update tensor at the given offset indices.

Traits: AlwaysSpeculatableImplTrait, AttrSizedOperandSegments

Interfaces: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface), TiedOpInterface

Effects: MemoryEffects::Effect{}

Operands:

Operand	Description
target	ranked tensor of any type values
target_dims	variadic of index
start_indices	variadic of index
update	ranked tensor of any type values
update_dims	variadic of index

Results:

Result	Description
result	ranked tensor of any type values

Utility ops

iree_input.align (Input::AlignOp)

Aligns up to a power-of-two alignment if required

Syntax:

operation ::= `iree_input.align` $value `,` $alignment attr-dict `:` type($result)

Aligns |value| up to the given power-of-two |alignment| if required.

Traits: AlwaysSpeculatableImplTrait, SameOperandsAndResultType

Interfaces: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Operand	Description
value	signless-integer-like
alignment	signless-integer-like

Results:

Result	Description
result	signless-integer-like

iree_input.null (Input::NullOp)

A null value

Syntax:

operation ::= `iree_input.null` attr-dict `:` type($result)

Initializes reference and variant types with a null value.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Results:

Result	Description
result	any type

Workgroup dispatch ops

iree_input.dispatch.workgroup.count (Input::DispatchWorkgroupCountOp)

Returns the total workgroup count of the grid

Syntax:

operation ::= `iree_input.dispatch.workgroup.count` `[` $dimension `]` attr-dict `:` type($result)

The total number of workgroups along each dimension in the dispatch grid.

Corresponds to the NumWorkgroups SPIR-V built-in and the gridDim CUDA built-in variable, only in the iree dialect the number of dimensions is not restricted to 3 (XYZ).

%x = iree_input.dispatch.workgroup.count[0] : index
%y = iree_input.dispatch.workgroup.count[1] : index

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Attribute	MLIR Type	Description
dimension	::mlir::IntegerAttr	index attribute

Results:

Result	Description
result	index

iree_input.dispatch.workgroup.id (Input::DispatchWorkgroupIDOp)

Returns the index of the current workgroup in the grid

Syntax:

operation ::= `iree_input.dispatch.workgroup.id` `[` $dimension `]` attr-dict `:` type($result)

The global workgroup ID of the current workgroup in the range of [0, iree_input.dispatch.workgroup.count) along each dimension.

Corresponds to the WorkgroupId SPIR-V built-in and the blockIdx CUDA built-in variable, only in the iree dialect the number of dimensions is not restricted to 3 (XYZ).

%x = iree_input.dispatch.workgroup.id[0] : index
%y = iree_input.dispatch.workgroup.id[1] : index

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Attribute	MLIR Type	Description
dimension	::mlir::IntegerAttr	index attribute

Results:

Result	Description
result	index

iree_input.dispatch.workgroup.size (Input::DispatchWorkgroupSizeOp)

Returns the size of each workgroup in invocations

Syntax:

operation ::= `iree_input.dispatch.workgroup.size` `[` $dimension `]` attr-dict `:` type($result)

The number of local invocations within the current workgroup along each dimension. Depending on backend this may map to the SIMT thread count or inner loop nest parameters.

Workgroup sizes are not determined at the iree dialect level as they are dependent on the target backend determined when lowering into the HAL. It's still possible to use the symbolic workgroup size inside of dispatch executables as a placeholder for the resolved value once in the HAL.

Corresponds to the WorkgroupSize SPIR-V built-in and the blockDim CUDA built-in variable, only in the iree dialect the number of dimensions is not restricted to 3 (XYZ).

%x = iree_input.dispatch.workgroup.size[0] : index
%y = iree_input.dispatch.workgroup.size[1] : index

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Attribute	MLIR Type	Description
dimension	::mlir::IntegerAttr	index attribute

Results:

Result	Description
result	index

Attributes

DescriptorSetBindingAttr

descriptor set binding specification

Syntax:

#iree_input.descriptor_set.binding<
  int64_t,   # ordinal
  DescriptorType,   # type
  std::optional<DescriptorFlags>   # flags
>

Parameters:

Parameter	C++ type	Description
ordinal	int64_t
type	DescriptorType
flags	std::optional<DescriptorFlags>

DescriptorSetLayoutAttr

descriptor set layout specification

Syntax:

#iree_input.descriptor_set.layout<
  int64_t,   # ordinal
  ::llvm::ArrayRef<DescriptorSetBindingAttr>,   # bindings
  std::optional<DescriptorSetLayoutFlags>   # flags
>

Parameters:

Parameter	C++ type	Description
ordinal	int64_t
bindings	::llvm::ArrayRef<DescriptorSetBindingAttr>
flags	std::optional<DescriptorSetLayoutFlags>

DescriptorTypeAttr

valid DescriptorType

Syntax:

#iree_input.descriptor_type<
  ::mlir::iree_compiler::IREE::Input::DescriptorType   # value
>

Enum cases: * uniform_buffer (UniformBuffer) * storage_buffer (StorageBuffer)

Parameters:

Parameter	C++ type	Description
value	::mlir::iree_compiler::IREE::Input::DescriptorType	an enum of type DescriptorType

DeviceTargetAttr

generic device target specification

Parameters:

Parameter	C++ type	Description
deviceID	StringAttr
configuration	DictionaryAttr

ExecutableObjectAttr

executable object reference

Parameters:

Parameter	C++ type	Description
path	StringAttr
data	DenseIntElementsAttr

ExecutableObjectsAttr

target-specific object file references

Parameters:

Parameter	C++ type	Description
targets	ArrayAttr
targetObjects	ArrayAttr

ExecutableTargetAttr

generic executable target specification

Parameters:

Parameter	C++ type	Description
backend	StringAttr
format	StringAttr
configuration	DictionaryAttr

PipelineLayoutAttr

executable entry point layout specification

Syntax:

#iree_input.pipeline.layout<
  int64_t,   # pushConstants
  ::llvm::ArrayRef<DescriptorSetLayoutAttr>   # setLayouts
>

Parameters:

Parameter	C++ type	Description
pushConstants	int64_t
setLayouts	::llvm::ArrayRef<DescriptorSetLayoutAttr>

Type constraints

list

A mutable, resizable list of some type.

Types

BufferType

Buffer is an untyped bag of bits with no shape or dtype

Syntax: !iree_input.buffer

Buffers represent an untyped bag of bits that can be reinterpreted depending on a use case using buffer_view operation. Buffers can be used for packing multiple tensors into the same underlying storage. It is left to higher level code to decide how exactly tensors layed out in the buffer.

BufferViewType

View into a buffer, with runtime shape and element type

Syntax: !iree_input.buffer_view

BufferViews represent views onto backing IREE runtime Buffer objects, adding runtime shape and element type parameters to the backing buffer. BufferViews are typically accepted and returned at boundaries with external code.

In the runtime and lower level compiler, BufferView's are fully modeled; however, as boundary types, not all features are exposed publicly. Since within compiled tensor programs, it is typical to operate in terms of fully typed tensors, the primary mechanism for getting or using a BufferView at the high level is by casting to/from a tensor. It is left to higher level code to ensure that aliasing rules are enforced at such boundaries.

ByteBufferType

a reference counted byte buffer

Syntax: !iree_input.byte_buffer

A reference counted byte buffer that models a pointer, offset, and length.

ListType

A one dimensional list of runtime values

Represents a list of arbitrary type. Primitive types can be expected to be efficiently stored in an unboxed form. Reference types and variants are permitted.

Lists can either be homogenous, with a fixed element type, or heterogenous by parameterizing them with a VariantType.

Parameters:

Parameter	C++ type	Description
elementType	::mlir::Type	A type suitable as an element type of a container

PtrType

Pointer to a concrete type

Parameters:

Parameter	C++ type	Description
targetType	::mlir::Type	A type suitable as a target type of a pointer

VariantType

Represents any legal or reference type in the IREE runtime

Syntax: !iree_input.variant

The variant type is typically used to parameterize container types that can contain any legal primitive, reference or null in the IREE type system.