'iree_input' Dialectlink

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

Operationslink

Buffer and buffer view opslink

`iree_input.buffer.subspan` (Input::BufferSubspanOp)link

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:link

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

Results:link

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

`iree_input.buffer_view.create` (Input::BufferViewCreateOp)link

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:link

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:link

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

`iree_input.buffer_view.dim` (Input::BufferViewDimOp)link

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:link

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

Operands:link

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

Results:link

Result	Description
`result`	index

`iree_input.buffer_view.rank` (Input::BufferViewRankOp)link

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:link

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

Results:link

Result	Description
`result`	index

Byte buffer opslink

`iree_input.byte_buffer.constant` (Input::ByteBufferConstantOp)link

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:link

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:link

Result	Description
`result`	a reference counted byte buffer

Compiler hint opslink

`iree_input.optimization_barrier` (Input::OptimizationBarrierOp)link

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:link

Operand	Description
`operands`	variadic of any type

Results:link

Result	Description
`results`	variadic of any type

Dispatch opslink

`iree_input.dispatch` (Input::DispatchOp)link

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:link

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

Operands:link

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

Results:link

Result	Description
`results`	variadic of any type

Executable source opslink

`iree_input.executable.export` (Input::ExecutableExportOp)link

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:link

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)link

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)link

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:link

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 opslink

`iree_input.global.address` (Input::GlobalAddressOp)link

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:link

Attribute	MLIR Type	Description
`global`	FlatSymbolRefAttr	symbol reference attribute

Results:link

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)link

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:link

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

Results:link

Result	Description
`result`	any type

`iree_input.global.load` (Input::GlobalLoadOp)link

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:link

Attribute	MLIR Type	Description
`global`	FlatSymbolRefAttr	symbol reference attribute

Results:link

Result	Description
`result`	any type

`iree_input.global` (Input::GlobalOp)link

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:link

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)link

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:link

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)link

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:link

Attribute	MLIR Type	Description
`global`	FlatSymbolRefAttr	symbol reference attribute

Operands:link

Operand	Description
`value`	any type

Mutable list opslink

`iree_input.list.create` (Input::ListCreateOp)link

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:link

Operand	Description
`initial_capacity`	index

Results:link

Result	Description
`result`	list

`iree_input.list.get` (Input::ListGetOp)link

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:link

Operand	Description
`list`	list
`index`	index

Results:link

Result	Description
`result`	any type

`iree_input.list.resize` (Input::ListResizeOp)link

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:link

Operand	Description
`list`	list
`new_size`	index

`iree_input.list.set` (Input::ListSetOp)link

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:link

Operand	Description
`list`	list
`index`	index
`value`	any type

`iree_input.list.size` (Input::ListSizeOp)link

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:link

Operand	Description
`list`	list

Results:link

Result	Description
`result`	index

Pseudo ops for conversion supportlink

`iree_input.tensor.export` (Input::TensorExportOp)link

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:link

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

Results:link

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)link

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:link

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:link

Result	Description
`target`	ranked tensor of any type values

Tensor opslink

`iree_input.tensor.bitcast` (Input::TensorBitCastOp)link

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:link

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

Results:link

Result	Description
`result`	ranked tensor of any type values

`iree_input.tensor.clone` (Input::TensorCloneOp)link

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:link

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

Results:link

Result	Description
`result`	ranked tensor of any type values

`iree_input.tensor.load` (Input::TensorLoadOp)link

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:link

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

Results:link

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)link

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:link

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

Results:link

Result	Description
`result`	ranked tensor of any type values

`iree_input.tensor.slice` (Input::TensorSliceOp)link

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:link

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:link

Result	Description
`result`	ranked tensor of any type values

`iree_input.tensor.splat` (Input::TensorSplatOp)link

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:link

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

Results:link

Result	Description
`result`	ranked tensor of any type values

`iree_input.tensor.store` (Input::TensorStoreOp)link

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:link

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:link

Result	Description
`result`	ranked tensor of any type values

`iree_input.tensor.trace` (Input::TensorTraceOp)link

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:link

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

Operands:link

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

`iree_input.tensor.update` (Input::TensorUpdateOp)link

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:link

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:link

Result	Description
`result`	ranked tensor of any type values

Utility opslink

`iree_input.align` (Input::AlignOp)link

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:link

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

Results:link

Result	Description
`result`	signless-integer-like

`iree_input.null` (Input::NullOp)link

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:link

Result	Description
`result`	any type

Workgroup dispatch opslink

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

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:link

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

Results:link

Result	Description
`result`	index

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

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:link

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

Results:link

Result	Description
`result`	index

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

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:link

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

Results:link

Result	Description
`result`	index

Attributeslink

DescriptorTypeAttrlink

valid DescriptorType

Syntax:

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

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

Parameters:link

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

DeviceTargetAttrlink

generic device target specification

Parameters:link

Parameter	C++ type	Description
deviceID	`StringAttr`
configuration	`DictionaryAttr`

ExecutableObjectAttrlink

executable object reference

Parameters:link

Parameter	C++ type	Description
path	`StringAttr`
data	`DenseIntElementsAttr`

ExecutableObjectsAttrlink

target-specific object file references

Parameters:link

Parameter	C++ type	Description
targets	`ArrayAttr`
targetObjects	`ArrayAttr`

ExecutableTargetAttrlink

generic executable target specification

Parameters:link

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

PipelineBindingAttrlink

descriptor set binding specification

Syntax:

#iree_input.pipeline.binding<
  DescriptorType,   # type
  std::optional<DescriptorFlags>   # flags
>

Parameters:link

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

PipelineLayoutAttrlink

executable entry point layout specification

Syntax:

#iree_input.pipeline.layout<
  ::llvm::ArrayRef<PipelineBindingAttr>,   # bindings
  int64_t,   # constants
  std::optional<PipelineLayoutFlags>   # flags
>

Parameters:link

Parameter	C++ type	Description
bindings	`::llvm::ArrayRef<PipelineBindingAttr>`
constants	`int64_t`	0
flags	`std::optional<PipelineLayoutFlags>`

Type constraintslink

listlink

A mutable, resizable list of some type.

Typeslink

BufferTypelink

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.

BufferViewTypelink

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.

ByteBufferTypelink

a reference counted byte buffer

Syntax: !iree_input.byte_buffer

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

ListTypelink

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:link

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

PtrTypelink

Pointer to a concrete type

Parameters:link

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

VariantTypelink

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.

Enumslink

DescriptorFlagslink

valid Descriptor flags

Cases:link

Symbol	Value	String
None	`0`	None
ReadOnly	`1`	ReadOnly

PipelineLayoutFlagslink

valid PipelineLayout flags

Cases:link

Symbol	Value	String
None	`0`	None
Indirect	`1`	Indirect

DescriptorTypelink

valid DescriptorType

Cases:link

Symbol	Value	String
UniformBuffer	`6`	uniform_buffer
StorageBuffer	`7`	storage_buffer

'iree_input' Dialectlink

Operationslink

Buffer and buffer view opslink

iree_input.buffer.subspan (Input::BufferSubspanOp)link

Operands:link

Results:link

iree_input.buffer_view.create (Input::BufferViewCreateOp)link

Operands:link

Results:link

iree_input.buffer_view.dim (Input::BufferViewDimOp)link

Attributes:link

Operands:link

Results:link

iree_input.buffer_view.rank (Input::BufferViewRankOp)link

Operands:link

Results:link

Byte buffer opslink

iree_input.byte_buffer.constant (Input::ByteBufferConstantOp)link

Attributes:link

Results:link

Compiler hint opslink

iree_input.optimization_barrier (Input::OptimizationBarrierOp)link

Operands:link

Results:link

Dispatch opslink

iree_input.dispatch (Input::DispatchOp)link

Attributes:link

Operands:link

Results:link

Executable source opslink

iree_input.executable.export (Input::ExecutableExportOp)link

Attributes:link

iree_input.executable.source_end (Input::ExecutableSourceEndOp)link

iree_input.executable.source (Input::ExecutableSourceOp)link

Attributes:link

Global variable opslink

iree_input.global.address (Input::GlobalAddressOp)link

Attributes:link

Results:link

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

Operands:link

Results:link

iree_input.global.load (Input::GlobalLoadOp)link

Attributes:link

Results:link

iree_input.global (Input::GlobalOp)link

Attributes:link

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

Operands:link

iree_input.global.store (Input::GlobalStoreOp)link

Attributes:link

Operands:link

Mutable list opslink

iree_input.list.create (Input::ListCreateOp)link

Operands:link

Results:link

iree_input.list.get (Input::ListGetOp)link

Operands:link

Results:link

iree_input.list.resize (Input::ListResizeOp)link

Operands:link

iree_input.list.set (Input::ListSetOp)link

Operands:link

iree_input.list.size (Input::ListSizeOp)link

Operands:link

Results:link

Pseudo ops for conversion supportlink

iree_input.tensor.export (Input::TensorExportOp)link

Operands:link

Results:link

iree_input.tensor.import (Input::TensorImportOp)link

Operands:link

Results:link

Tensor opslink

iree_input.tensor.bitcast (Input::TensorBitCastOp)link

Operands:link

Results:link

iree_input.tensor.clone (Input::TensorCloneOp)link

Operands:link

Results:link

`iree_input.buffer.subspan` (Input::BufferSubspanOp)link

`iree_input.buffer_view.create` (Input::BufferViewCreateOp)link

`iree_input.buffer_view.dim` (Input::BufferViewDimOp)link

`iree_input.buffer_view.rank` (Input::BufferViewRankOp)link

`iree_input.byte_buffer.constant` (Input::ByteBufferConstantOp)link

`iree_input.optimization_barrier` (Input::OptimizationBarrierOp)link

`iree_input.dispatch` (Input::DispatchOp)link

`iree_input.executable.export` (Input::ExecutableExportOp)link

`iree_input.executable.source_end` (Input::ExecutableSourceEndOp)link

`iree_input.executable.source` (Input::ExecutableSourceOp)link

`iree_input.global.address` (Input::GlobalAddressOp)link

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

`iree_input.global.load` (Input::GlobalLoadOp)link

`iree_input.global` (Input::GlobalOp)link

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

`iree_input.global.store` (Input::GlobalStoreOp)link

`iree_input.list.create` (Input::ListCreateOp)link

`iree_input.list.get` (Input::ListGetOp)link

`iree_input.list.resize` (Input::ListResizeOp)link

`iree_input.list.set` (Input::ListSetOp)link

`iree_input.list.size` (Input::ListSizeOp)link

`iree_input.tensor.export` (Input::TensorExportOp)link

`iree_input.tensor.import` (Input::TensorImportOp)link

`iree_input.tensor.bitcast` (Input::TensorBitCastOp)link

`iree_input.tensor.clone` (Input::TensorCloneOp)link

`iree_input.tensor.load` (Input::TensorLoadOp)link

`iree_input.tensor.reshape` (Input::TensorReshapeOp)link

`iree_input.tensor.slice` (Input::TensorSliceOp)link

`iree_input.tensor.splat` (Input::TensorSplatOp)link

`iree_input.tensor.store` (Input::TensorStoreOp)link

`iree_input.tensor.trace` (Input::TensorTraceOp)link

`iree_input.tensor.update` (Input::TensorUpdateOp)link

`iree_input.align` (Input::AlignOp)link

`iree_input.null` (Input::NullOp)link

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

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

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