Skip to content

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

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
AttributeMLIR TypeDescription
index::mlir::IntegerAttrindex 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
AttributeMLIR TypeDescription
name::mlir::StringAttrstring attribute
value::mlir::StringAttrstring attribute
alignment::mlir::IntegerAttrindex attribute
mime_type::mlir::StringAttrstring 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
AttributeMLIR TypeDescription
entry_point::mlir::SymbolRefAttrsymbol reference attribute
tied_operands::mlir::ArrayAttr64-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
AttributeMLIR TypeDescription
sym_visibility::mlir::StringAttrstring attribute
sym_name::mlir::StringAttrstring attribute
ordinal::mlir::IntegerAttrsize_t
layout::mlir::iree_compiler::IREE::Input::PipelineLayoutAttrexecutable entry point layout specification
workgroup_size::mlir::ArrayAttrindex array attribute
subgroup_size::mlir::IntegerAttrsize_t
workgroup_local_memory::mlir::IntegerAttrindex 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
AttributeMLIR TypeDescription
sym_visibility::mlir::StringAttrstring attribute
sym_name::mlir::StringAttrstring attribute
objects::mlir::iree_compiler::IREE::Input::ExecutableObjectsAttrtarget-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
AttributeMLIR TypeDescription
globalFlatSymbolRefAttrsymbol 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
AttributeMLIR TypeDescription
globalFlatSymbolRefAttrsymbol 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
AttributeMLIR TypeDescription
sym_visibility::mlir::StringAttrstring attribute
sym_name::mlir::StringAttrstring attribute
type::mlir::TypeAttrany type attribute
is_mutable::mlir::UnitAttrunit attribute
initializer::mlir::FlatSymbolRefAttrflat symbol reference attribute
initial_value::mlir::TypedAttrTypedAttr 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
AttributeMLIR TypeDescription
globalFlatSymbolRefAttrsymbol 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
AttributeMLIR TypeDescription
key::mlir::StringAttrstring 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
AttributeMLIR TypeDescription
dimension::mlir::IntegerAttrindex 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
AttributeMLIR TypeDescription
dimension::mlir::IntegerAttrindex 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
AttributeMLIR TypeDescription
dimension::mlir::IntegerAttrindex 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