Skip to content

'hal' Dialectlink

A dialect representing operations against the IREE HAL.

This can be thought of as a Vulkan-like model with all of the graphics bits chopped out.

The type set is limited to those that can be represented in the IREE HAL design: buffers and views, synchronization primitives like semaphores, and command buffers. The intent is that if a device could implement the HAL interface the sequencer ops could run on that device, such as being able to run on a GPU via indirect command buffers.

Though this is mostly a 1:1 mapping to the iree::hal API there are some methods omitted as they are not likely to be needed in IR. It's assumed that either sequencer interfaces will encapsulate the logic (such as device resolution) or that certain features are unsafe to expose to user-defined input.

Operationslink

Allocator opslink

Ops for !hal.allocator / iree_hal_allocator_t.

hal.allocator.allocate (HAL::AllocatorAllocateOp)link

Empty buffer allocation operation

Syntax:

operation ::= `hal.allocator.allocate` `<` $allocator `:` type($allocator) `>`
              `affinity` `(` $queue_affinity `)`
              `type` `(` $memory_types `)`
              `usage` `(` $buffer_usage `)`
              `:` custom<SizeAwareType>(type($result), $result_size)
              attr-dict-with-keyword

Allocates a buffer of the given size from the allocator. The size of the buffer returned may be larger than the requested size if the allocator has specific alignment requirements or minimum allocation sizes.

Interfaces: OpAsmOpInterface, SizeAwareOpInterface

Attributes:link
AttributeMLIR TypeDescription
memory_typesmlir::iree_compiler::IREE::HAL::MemoryTypeBitfieldAttrvalid MemoryType
buffer_usagemlir::iree_compiler::IREE::HAL::BufferUsageBitfieldAttrvalid BufferUsage
Operands:link
Operand Description
allocator allocator
queue_affinity 64-bit signless integer
result_size index
Results:link
Result Description
result buffer

hal.allocator.import (HAL::AllocatorImportOp)link

Allocator-supported host buffer import operation

Syntax:

operation ::= `hal.allocator.import` `<` $allocator `:` type($allocator) `>`
              `source` `(` $source `:` type($source) `)` `` `[` $offset `,` $length `]`
              `affinity` `(` $queue_affinity `)`
              `type` `(` $memory_types `)`
              `usage` `(` $buffer_usage `)`
              `:` type($did_import) `,` type($result)
              attr-dict-with-keyword

Tries importing host memory backed by the given byte buffer into a device accessible !hal.buffer. The returned buffer may be host-only and not directly usable on devices. If the mapping cannot be completed (such as trying to map the host memory as device-local on devices with discrete memory) then did_import will indicate that the returned buffer is null.

Interfaces: OpAsmOpInterface, SizeAwareOpInterface

Attributes:link
AttributeMLIR TypeDescription
memory_typesmlir::iree_compiler::IREE::HAL::MemoryTypeBitfieldAttrvalid MemoryType
buffer_usagemlir::iree_compiler::IREE::HAL::BufferUsageBitfieldAttrvalid BufferUsage
Operands:link
Operand Description
allocator allocator
queue_affinity 64-bit signless integer
source a reference counted byte buffer
offset index
length index
Results:link
Result Description
did_import 1-bit signless integer
result buffer

Buffer opslink

Ops for !hal.buffer / iree_hal_buffer_t.

hal.buffer.assert (HAL::BufferAssertOp)link

Buffer compatibility assertion

Syntax:

operation ::= `hal.buffer.assert` `<` $buffer `:` type($buffer) `>`
              `message` `(` $message `)`
              `allocator` `(` $allocator `:` type($allocator) `)`
              `minimum_length` `(` $minimum_length `)`
              `type` `(` $memory_types `)`
              `usage` `(` $buffer_usage `)`
              attr-dict-with-keyword

Asserts that the buffer is compatible with the given allocator and usage. Program execution will abort as if std.assert had been used.

This only checks that the buffer can be used and not that it matches the given parameters exactly. Buffers may be from other allocators so long as the allocators are compatible (devices can address each other's memory), the type and usage contain all the requested bits (having more bits is ok), and the length is at least the requested minimum (as padding may be ignored).

Attributes:link
AttributeMLIR TypeDescription
message::mlir::StringAttrstring attribute
memory_typesmlir::iree_compiler::IREE::HAL::MemoryTypeBitfieldAttrvalid MemoryType
buffer_usagemlir::iree_compiler::IREE::HAL::BufferUsageBitfieldAttrvalid BufferUsage
Operands:link
Operand Description
buffer buffer
allocator allocator
minimum_length index

hal.buffer.length (HAL::BufferLengthOp)link

Buffer byte length accessor

Syntax:

operation ::= `hal.buffer.length` `<` $buffer `:` type($buffer) `>`
              `:` type($result)
              attr-dict-with-keyword

Returns the allocated size of a buffer in bytes. May be less than the underlying buffer allocation if this is a subspan or view into another buffer.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface), OpAsmOpInterface

Effects: MemoryEffects::Effect{}

Operands:link
Operand Description
buffer buffer
Results:link
Result Description
result index

hal.buffer.load (HAL::BufferLoadOp)link

Buffer element load operation

Syntax:

operation ::= `hal.buffer.load` `<` $source_buffer `:` type($source_buffer) `>`
              `` `[` $source_offset `]`
              `:` type($result)
              attr-dict-with-keyword

Loads a value from a buffer by mapping it.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:link
Operand Description
source_buffer buffer
source_offset index
Results:link
Result Description
result index or signless integer or floating-point or complex-type or vector of any type values

hal.buffer.store (HAL::BufferStoreOp)link

Buffer element store operation

Syntax:

operation ::= `hal.buffer.store` `<` $target_buffer `:` type($target_buffer) `>`
              `` `[` $target_offset `]`
              `value` `(` $value `:` type($value) `)`
              attr-dict-with-keyword

Stores a value into a buffer by mapping it.

Operands:link
Operand Description
value index or signless integer or floating-point or complex-type or vector of any type values
target_buffer buffer
target_offset index

hal.buffer.subspan (HAL::BufferSubspanOp)link

Buffer subspan operation

Syntax:

operation ::= `hal.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, NoMemoryEffect (MemoryEffectOpInterface), OpAsmOpInterface, SizeAwareOpInterface

Effects: MemoryEffects::Effect{}

Operands:link
Operand Description
source_buffer buffer
source_offset index
length index
Results:link
Result Description
result buffer

Buffer view opslink

Ops for !hal.buffer_view / iree_hal_buffer_view_t.

hal.buffer_view.assert (HAL::BufferViewAssertOp)link

Buffer view contents assertion

Syntax:

operation ::= `hal.buffer_view.assert` `<` $buffer_view `:` type($buffer_view) `>`
              `message` `(` $message `)`
              `shape` `(` `[` $shape `]` `)`
              `type` `(` $element_type `)`
              `encoding` `(` $encoding_type `)`
              attr-dict-with-keyword

Asserts that the buffer view contains a data compatible tensor with the given encoding. Program execution will abort as if std.assert had been used.

Attributes:link
AttributeMLIR TypeDescription
message::mlir::StringAttrstring attribute
Operands:link
Operand Description
buffer_view buffer_view
element_type 32-bit signless integer
encoding_type 32-bit signless integer
shape variadic of index

hal.buffer_view.buffer (HAL::BufferViewBufferOp)link

Buffer view buffer accessor

Syntax:

operation ::= `hal.buffer_view.buffer` `<` $buffer_view `:` type($buffer_view) `>`
              `:` type($result)
              attr-dict-with-keyword

Returns the buffer backing this view's contents.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface), OpAsmOpInterface

Effects: MemoryEffects::Effect{}

Operands:link
Operand Description
buffer_view buffer_view
Results:link
Result Description
result buffer

hal.buffer_view.create (HAL::BufferViewCreateOp)link

Buffer view reference initializer

Syntax:

operation ::= `hal.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, NoMemoryEffect (MemoryEffectOpInterface), OpAsmOpInterface

Effects: MemoryEffects::Effect{}

Operands:link
Operand Description
source_buffer buffer
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 buffer_view

hal.buffer_view.dim (HAL::BufferViewDimOp)link

Buffer view dimension value query

Syntax:

operation ::= `hal.buffer_view.dim` `<` $buffer_view `:` type($buffer_view) `>`
              `` `[` $index `]`
              `:` type($result)
              attr-dict-with-keyword

Returns the value of the given dimension.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:link
AttributeMLIR TypeDescription
index::mlir::IntegerAttrindex attribute
Operands:link
Operand Description
buffer_view buffer_view
Results:link
Result Description
result index

hal.buffer_view.element_type (HAL::BufferViewElementTypeOp)link

Buffer view element type query

Syntax:

operation ::= `hal.buffer_view.element_type` `<` $buffer_view `:` type($buffer_view) `>`
              `:` type($result)
              attr-dict-with-keyword

Returns the element type of the buffer view.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:link
Operand Description
buffer_view buffer_view
Results:link
Result Description
result 32-bit signless integer

hal.buffer_view.encoding_type (HAL::BufferViewEncodingTypeOp)link

Buffer view encoding type query

Syntax:

operation ::= `hal.buffer_view.encoding_type` `<` $buffer_view `:` type($buffer_view) `>`
              `:` type($result)
              attr-dict-with-keyword

Returns the encoding type of the buffer view.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:link
Operand Description
buffer_view buffer_view
Results:link
Result Description
result 32-bit signless integer

hal.buffer_view.rank (HAL::BufferViewRankOp)link

Buffer view rank query

Syntax:

operation ::= `hal.buffer_view.rank` `<` $buffer_view `:` type($buffer_view) `>`
              `:` type($result)
              attr-dict-with-keyword

Returns the rank of the buffer view.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:link
Operand Description
buffer_view buffer_view
Results:link
Result Description
result index

hal.buffer_view.trace (HAL::BufferViewTraceOp)link

Trace value(s) operation

Syntax:

operation ::= `hal.buffer_view.trace` $key `=`
              $operands `:` type($operands)
              attr-dict-with-keyword

Traces out to a runtime trace sink (console, log file, etc) the given buffer views and titles them with the given key. The key is informational only and useful for titling/marking specific sets of buffers for easier searching.

Attributes:link
AttributeMLIR TypeDescription
key::mlir::StringAttrstring attribute
Operands:link
Operand Description
operands variadic of buffer_view

hal.element_type (HAL::ElementTypeOp)link

An iree_hal_element_type_t for the given MLIR type

Syntax:

operation ::= `hal.element_type` `<` $type `>`
              attr-dict
              `:` type($result)

Maps an MLIR type to a runtime iree_hal_element_type_t value for all types that are convertable.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface), OpAsmOpInterface

Effects: MemoryEffects::Effect{}

Attributes:link
AttributeMLIR TypeDescription
type::mlir::TypeAttrany type attribute
Results:link
Result Description
result 32-bit signless integer

hal.encoding_type (HAL::EncodingTypeOp)link

An iree_hal_encoding_type_t for the given MLIR encoding

Syntax:

operation ::= `hal.encoding_type` `<` ($encoding^):( `` `dense_row_major`)? `>`
              attr-dict
              `:` type($result)

Maps an MLIR encoding to a runtime iree_hal_encoding_type_t value for all encodings that are convertable.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface), OpAsmOpInterface

Effects: MemoryEffects::Effect{}

Attributes:link
AttributeMLIR TypeDescription
encoding::mlir::Attributeany attribute
Results:link
Result Description
result 32-bit signless integer

Channel opslink

Ops for !hal.channel / iree_hal_channel_t.

hal.channel.create (HAL::ChannelCreateOp)link

Creates a new channel for collective communication

Syntax:

operation ::= `hal.channel.create` `device` `(` $device `:` type($device) `)`
              `affinity` `(` $queue_affinity `)`
              `flags` `(` $flags `)`
              `id` `(` $id `)`
              `group` `(` $group `)`
              `rank` `(` $rank `)`
              `count` `(` $count `)`
              `:` type($result)
              attr-dict-with-keyword

Returns a new channel with the given rank associated with the given device queue. Collective operations using this channel must only be submitted on compatible queues.

The group and ID are optional and may be null. A rank or count of -1 can be used to indicate a default inherited from the environment or device configuration.

Interfaces: OpAsmOpInterface

Attributes:link
AttributeMLIR TypeDescription
flags::mlir::IntegerAttr32-bit signless integer attribute
Operands:link
Operand Description
device device
queue_affinity 64-bit signless integer
id a reference counted byte buffer
group a reference counted byte buffer
rank 32-bit signless integer
count 32-bit signless integer
Results:link
Result Description
result collective.channel

hal.channel.rank_and_count (HAL::ChannelRankAndCountOp)link

Returns the rank of the local participant in the group

Syntax:

operation ::= `hal.channel.rank_and_count` `<` $channel `:` type($channel) `>`
              `:` type($rank) `,` type($count)
              attr-dict-with-keyword

Returns the rank the channel represents as a participant in a collective group in [0, count) and the total participant count.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface), OpAsmOpInterface

Effects: MemoryEffects::Effect{}

Operands:link
Operand Description
channel collective.channel
Results:link
Result Description
rank 32-bit signless integer
count 32-bit signless integer

hal.channel.split (HAL::ChannelSplitOp)link

Splits a collective communication channel

Syntax:

operation ::= `hal.channel.split` `<` $channel `:` type($channel) `>`
              `color` `(` $color `)`
              `key` `(` $key `)`
              `flags` `(` $flags `)`
              `:` type($result)
              attr-dict-with-keyword

Partitions the group associated with the given channel into disjoint subgroups for each unique value of color. Each new subgroup contains all participants of the same color and within each subgroup the key argument is used to define the rank order. When multiple participants in a group use the same key the tie will be broken using their rank in the parent group. A color of -1 indicates that the rank does not participate in any subgroup and will return a null channel.

Interfaces: OpAsmOpInterface

Attributes:link
AttributeMLIR TypeDescription
flags::mlir::IntegerAttr32-bit signless integer attribute
Operands:link
Operand Description
channel collective.channel
color 32-bit signless integer
key 32-bit signless integer
Results:link
Result Description
result collective.channel

Command buffer opslink

Ops for !hal.command_buffer / iree_hal_command_buffer_t.

hal.command_buffer.begin_debug_group (HAL::CommandBufferBeginDebugGroupOp)link

Pushes a command buffer debug group label

Syntax:

operation ::= `hal.command_buffer.begin_debug_group` `<` $command_buffer `:` type($command_buffer) `>`
              `label` `(` $label `)`
              attr-dict-with-keyword

Pushes a new debug group with the given label. All commands between this and a mandatory matching call to hal.command_buffer.end_debug_group will be grouped together with the given label.

Attributes:link
AttributeMLIR TypeDescription
label::mlir::StringAttrstring attribute
Operands:link
Operand Description
command_buffer command_buffer

hal.command_buffer.collective (HAL::CommandBufferCollectiveOp)link

Command buffer collective dispatch recording operation

Syntax:

operation ::= `hal.command_buffer.collective` `<` $command_buffer `:` type($command_buffer) `>`
              `channel` `(` $channel `:` type($channel) `)`
              `op` `(` $op `)`
              (`param` `(` $param^ `:` type($param) `)`)?
              (`send` `(` $send_buffer^ `:` type($send_buffer) `)`
              `` `[` $send_offset `,` $send_length `]`)?
              (`recv` `(` $recv_buffer^ `:` type($recv_buffer) `)`
              `` `[` $recv_offset `,` $recv_length `]`)?
              `count` `(` $element_count `)`
              attr-dict-with-keyword

Dispatches a collective operation defined by op using the given buffers.

Traits: AttrSizedOperandSegments

Attributes:link
AttributeMLIR TypeDescription
op::mlir::iree_compiler::IREE::HAL::CollectiveAttrcollective operation and specification
Operands:link
Operand Description
command_buffer command_buffer
channel collective.channel
element_count index
param 32-bit signless integer
send_buffer buffer
send_offset index
send_length index
recv_buffer buffer
recv_offset index
recv_length index

hal.command_buffer.copy_buffer (HAL::CommandBufferCopyBufferOp)link

Command buffer buffer copy recording operation

Syntax:

operation ::= `hal.command_buffer.copy_buffer` `<` $command_buffer `:` type($command_buffer) `>`
              `source` `(` $source_buffer `:` type($source_buffer) `)`
              `` `[` $source_offset `]`
              `target` `(` $target_buffer `:` type($target_buffer) `)`
              `` `[` $target_offset `]`
              `length` `(` $length `)`
              attr-dict-with-keyword

Copies a range of one buffer to another.

Operands:link
Operand Description
command_buffer command_buffer
source_buffer buffer
source_offset index
target_buffer buffer
target_offset index
length index

hal.command_buffer.create (HAL::CommandBufferCreateOp)link

Command buffer allocation operation

Syntax:

operation ::= `hal.command_buffer.create` `device` `(` $device `:` type($device) `)`
              `mode` `(` $modes `)`
              `categories` `(` $command_categories `)`
              (`bindings` `(` $binding_capacity^ `)`)?
              `:` type($result)
              attr-dict-with-keyword

Returns a command buffer from the device pool ready to begin recording.

Interfaces: OpAsmOpInterface

Attributes:link
AttributeMLIR TypeDescription
modesmlir::iree_compiler::IREE::HAL::CommandBufferModeBitfieldAttrvalid CommandBufferMode
command_categoriesmlir::iree_compiler::IREE::HAL::CommandCategoryBitfieldAttrvalid CommandCategory
Operands:link
Operand Description
device device
binding_capacity index
Results:link
Result Description
result command_buffer

hal.command_buffer.device (HAL::CommandBufferDeviceOp)link

Command buffer device query operation

Syntax:

operation ::= `hal.command_buffer.device` `<` $command_buffer `:` type($command_buffer) `>`
              `:` type($device)
              attr-dict-with-keyword

Used during conversion to access the device used to create a command buffer.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:link
Operand Description
command_buffer command_buffer
Results:link
Result Description
device device

hal.command_buffer.dispatch.indirect (HAL::CommandBufferDispatchIndirectOp)link

Command buffer indirect dispatch recording operation

Syntax:

operation ::= `hal.command_buffer.dispatch.indirect` `<` $command_buffer `:` type($command_buffer) `>`
              `target` `(` $executable `:` type($executable) `)`
              `` `[` $entry_point `]`
              `workgroups` `(` $workgroups_buffer `:` type($workgroups_buffer) `)`
              `` `[` $workgroups_offset `]`
              attr-dict-with-keyword

Dispatches an execution request with the dispatch parameters loaded from the given buffer.

Operands:link
Operand Description
command_buffer command_buffer
executable executable
entry_point index
workgroups_buffer buffer
workgroups_offset index

hal.command_buffer.dispatch (HAL::CommandBufferDispatchOp)link

Command buffer dispatch recording operation

Syntax:

operation ::= `hal.command_buffer.dispatch` `<` $command_buffer `:` type($command_buffer) `>`
              `target` `(` $executable `:` type($executable) `)`
              `` `[` $entry_point `]`
              `workgroups` `(` `[`
              $workgroup_x `,`
              $workgroup_y `,`
              $workgroup_z
              `]` `)`
              attr-dict-with-keyword

Dispatches an execution request.

Operands:link
Operand Description
command_buffer command_buffer
executable executable
entry_point index
workgroup_x index
workgroup_y index
workgroup_z index

hal.command_buffer.end_debug_group (HAL::CommandBufferEndDebugGroupOp)link

Pops a command buffer debug group label

Syntax:

operation ::= `hal.command_buffer.end_debug_group` `<` $command_buffer `:` type($command_buffer) `>`
              attr-dict-with-keyword

Pops a debug group from the stack.

Operands:link
Operand Description
command_buffer command_buffer

hal.command_buffer.execution_barrier (HAL::CommandBufferExecutionBarrierOp)link

Command buffer execution barrier recording operation

Syntax:

operation ::= `hal.command_buffer.execution_barrier` `<` $command_buffer `:` type($command_buffer) `>`
              `source` `(` $source_stage_mask `)`
              `target` `(` $target_stage_mask `)`
              `flags` `(` $flags `)`
              attr-dict-with-keyword

Defines an execution dependency between all commands recorded before the barrier and all commands recorded after the barrier. Only the stages provided will be affected.

Attributes:link
AttributeMLIR TypeDescription
source_stage_maskmlir::iree_compiler::IREE::HAL::ExecutionStageBitfieldAttrvalid ExecutionStage
target_stage_maskmlir::iree_compiler::IREE::HAL::ExecutionStageBitfieldAttrvalid ExecutionStage
flagsmlir::iree_compiler::IREE::HAL::ExecutionBarrierFlagBitfieldAttrvalid ExecutionBarrierFlag
Operands:link
Operand Description
command_buffer command_buffer

hal.command_buffer.fill_buffer (HAL::CommandBufferFillBufferOp)link

Command buffer buffer fill recording operation

Syntax:

operation ::= `hal.command_buffer.fill_buffer` `<` $command_buffer `:` type($command_buffer) `>`
              `target` `(` $target_buffer `:` type($target_buffer) `)`
              `` `[` $target_offset `,` $length `]`
              `pattern` `(` $pattern `:` type($pattern) `)`
              attr-dict-with-keyword

Fills the target buffer with the given repeating value.

Operands:link
Operand Description
command_buffer command_buffer
target_buffer buffer
target_offset index
length index
pattern 8-bit signless integer or 16-bit signless integer or 32-bit signless integer

hal.command_buffer.finalize (HAL::CommandBufferFinalizeOp)link

Finalizes command buffer recording

Syntax:

operation ::= `hal.command_buffer.finalize` `<` $command_buffer `:` type($command_buffer) `>`
              attr-dict-with-keyword

Ends recording into the command buffer and prepares it for submission. No more commands may be recorded into the command buffer.

Operands:link
Operand Description
command_buffer command_buffer

hal.command_buffer.push_constants (HAL::CommandBufferPushConstantsOp)link

Command buffer push constants operation

Syntax:

operation ::= `hal.command_buffer.push_constants` `<` $command_buffer `:` type($command_buffer) `>`
              `layout` `(` $pipeline_layout `:` type($pipeline_layout) `)`
              `offset` `(` $offset `)`
              `values` `(` `[` $values `]` `)`
              `:` type($values)
              attr-dict-with-keyword

Pushes an inline set of constants that can be accessed by subsequent dispatches using a compatible pipeline layout.

Push constants are always 4-byte values and treated as opaque, meaning that they may be bit-casted floats, bit-packed booleans, etc.

Attributes:link
AttributeMLIR TypeDescription
offset::mlir::IntegerAttrindex attribute
Operands:link
Operand Description
command_buffer command_buffer
pipeline_layout pipeline_layout
values variadic of 32-bit signless integer

hal.command_buffer.push_descriptor_set (HAL::CommandBufferPushDescriptorSetOp)link

Command buffer descriptor set push binding operation

Syntax:

operation ::= `hal.command_buffer.push_descriptor_set` `<` $command_buffer `:` type($command_buffer) `>`
              `layout` `(` $pipeline_layout `:` type($pipeline_layout) `)`
              `` `[` $set `]`
              `bindings` `(` `[`
              custom<DescriptorSetBindings>($binding_ordinals,
              $binding_buffers,
              type($binding_buffers),
              $binding_offsets,
              $binding_lengths)
              `]` `)`
              attr-dict-with-keyword

Pushes an inline-defined descriptor set to the command buffer. The provided buffers may either be HAL buffers or indirect references into the command buffer binding table.

Traits: SameVariadicOperandSize

Operands:link
Operand Description
command_buffer command_buffer
pipeline_layout pipeline_layout
set index
binding_ordinals variadic of index
binding_buffers variadic of index or buffer
binding_offsets variadic of index
binding_lengths variadic of index

Descriptor set layout opslink

Ops for !hal.descriptor_set_layout / iree_hal_descriptor_set_layout_t.

hal.descriptor_set_layout.create (HAL::DescriptorSetLayoutCreateOp)link

Creates a descriptor set layout

Syntax:

operation ::= `hal.descriptor_set_layout.create` `device` `(` $device `:` type($device) `)`
              `flags` `(` $flags `)`
              `bindings` `(` $bindings `)`
              `:` type($result)
              attr-dict-with-keyword

Creates a descriptor set layout that defines the bindings used within a set. The same descriptor set layout may be shared with many different executable layouts and by doing so some runtime binding overhead when switching between executables that use the same set layouts can be reduced.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface), OpAsmOpInterface

Effects: MemoryEffects::Effect{}

Attributes:link
AttributeMLIR TypeDescription
flags::mlir::iree_compiler::IREE::HAL::DescriptorSetLayoutFlagsAttrvalid DescriptorSetLayout flags
bindings::mlir::ArrayAttrHAL descriptor set layout binding array attribute
Operands:link
Operand Description
device device
Results:link
Result Description
result descriptor_set_layout

Device management opslink

Device availability and selection support.

hal.devices.count (HAL::DevicesCountOp)link

Returns the number of available devices

Syntax:

operation ::= `hal.devices.count` attr-dict `:` type($result)

Returns the total number of available devices registered at runtime.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface), OpAsmOpInterface

Effects: MemoryEffects::Effect{}

Results:link
Result Description
result index

hal.devices.get (HAL::DevicesGetOp)link

Returns the device with the given index

Syntax:

operation ::= `hal.devices.get` $index attr-dict `:` type($result)

Returns the device with the given index in the [0, hal.devices.count) range. Devices may be lazily initialized upon first use.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface), OpAsmOpInterface

Effects: MemoryEffects::Effect{}

Operands:link
Operand Description
index index
Results:link
Result Description
result device

Device opslink

Ops for !hal.device / iree_hal_device_t.

hal.device.allocator (HAL::DeviceAllocatorOp)link

Device allocator accessor operation

Syntax:

operation ::= `hal.device.allocator` `<` $device `:` type($device) `>` `:` type($result) attr-dict-with-keyword

Returns the allocator that can be used to allocate buffers compatible with the device.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface), OpAsmOpInterface

Effects: MemoryEffects::Effect{}

Operands:link
Operand Description
device device
Results:link
Result Description
result allocator

hal.device.query (HAL::DeviceQueryOp)link

Returns a runtime configuration parameter from the device

Syntax:

operation ::= `hal.device.query` `<` $device `:` type($device) `>`
              `key` `(` $category `:` `` `:` $key `)`
              `:` type($ok) `,` type($value)
              (`=` $default_value^)?
              attr-dict-with-keyword

Queries a device configuration parameter with the given key. Returns a status indicating whether the pair was recognized/available and if it was the value converted to the specified type. Queries must return the same value for the lifetime of the module though may vary from run to run.

This is roughly equivalent to the sysconf linux syscall (https://man7.org/linux/man-pages/man3/sysconf.3.html) in that the exact set of keys available and their interpretation is target-dependent.

Users of the op must check the ok result before using the value as what set of keys is available may change over time. If in doubt: don't use this. Each key used adds additional versioning and testing complexity as runtime code path changes will explode combinatorially and should be treated with as much care as a binary file format change. Keys should be prefixed with ex. when experimental indicating that they are not expected to be present forever; all non-experimental keys should be vetted.

Well-known keys:

  • hal.device.id :: {some id pattern} Returns 1 if the device identifier matches the given pattern string.

  • hal.executable.format :: {some format pattern} Returns 1 if the given format is supported by the device loader.

  • hal.device :: concurrency The maximum concurrently executable submissions, mapping roughly to the queue count. The actual concurrency available may be less than this based on dynamic runtime parameters such as power/thermal modes, quota limits, or user choice.

  • hal.dispatch :: concurrency The maximum concurrently executable workgroups for a particular dispatch. The actual concurrency available may be less depending on device state.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:link
AttributeMLIR TypeDescription
category::mlir::StringAttrstring attribute
key::mlir::StringAttrstring attribute
default_value::mlir::TypedAttrTypedAttr instance
Operands:link
Operand Description
device device
Results:link
Result Description
ok 1-bit signless integer
value any type

hal.device.queue.alloca (HAL::DeviceQueueAllocaOp)link

Allocates a queue-ordered transient buffer

Syntax:

operation ::= `hal.device.queue.alloca` `<` $device `:` type($device) `>`
              `affinity` `(` $queue_affinity `)`
              `wait` `(` $wait_fence `)`
              `signal` `(` $signal_fence `)`
              `pool` `(` $pool `)`
              `type` `(` $memory_types `)`
              `usage` `(` $buffer_usage `)`
              `:` custom<SizeAwareType>(type($result), $result_size)
              attr-dict-with-keyword

Returns a queue-ordered transient buffer that will be available for use when the signal fence is reached. The allocation will not be made until the wait fence has been reached.

The size of the buffer returned may be larger than the requested size if the allocator has specific alignment requirements or minimum allocation sizes.

The buffer handle will remain live so long as there are retainers but the contents are undefined before the allocation signal fence has been signaled and after the deallocation wait fence has been reached.

Interfaces: OpAsmOpInterface, SizeAwareOpInterface

Attributes:link
AttributeMLIR TypeDescription
memory_typesmlir::iree_compiler::IREE::HAL::MemoryTypeBitfieldAttrvalid MemoryType
buffer_usagemlir::iree_compiler::IREE::HAL::BufferUsageBitfieldAttrvalid BufferUsage
Operands:link
Operand Description
device device
queue_affinity 64-bit signless integer
wait_fence fence
signal_fence fence
pool 64-bit signless integer
result_size index
Results:link
Result Description
result buffer

hal.device.queue.dealloca (HAL::DeviceQueueDeallocaOp)link

Deallocates a queue-ordered transient buffer

Syntax:

operation ::= `hal.device.queue.dealloca` `<` $device `:` type($device) `>`
              `affinity` `(` $queue_affinity `)`
              `wait` `(` $wait_fence `)`
              `signal` `(` $signal_fence `)`
              `buffer` `(` $buffer `:` type($buffer) `)`
              attr-dict-with-keyword

Deallocates a queue-ordered transient buffer. The deallocation will not be made until the wait fence has been reached and once the storage is available for reuse the signal fence will be signaled.

After deallocation the contents of the buffer may still be accessible but will have undefined contents as other operations reuse the memory.

Operands:link
Operand Description
device device
queue_affinity 64-bit signless integer
wait_fence fence
signal_fence fence
buffer buffer

hal.device.queue.execute.indirect (HAL::DeviceQueueExecuteIndirectOp)link

Enqueues command buffer execution

Syntax:

operation ::= `hal.device.queue.execute.indirect` `<` $device `:` type($device) `>`
              `affinity` `(` $queue_affinity `)`
              `wait` `(` $wait_fence `)`
              `signal` `(` $signal_fence `)`
              `commands` `(` $command_buffer `)`
              `bindings` `(` `[`
              custom<BindingTable>($binding_buffers,
              type($binding_buffers),
              $binding_offsets,
              $binding_lengths)
              `]` `)`
              attr-dict-with-keyword

Executes a command buffer on a device queue with the given binding table. No commands will execute until the wait fence has been reached and the signal fence will be signaled when all commands have completed.

Traits: SameVariadicOperandSize

Operands:link
Operand Description
device device
queue_affinity 64-bit signless integer
wait_fence fence
signal_fence fence
command_buffer command_buffer
binding_buffers variadic of buffer
binding_offsets variadic of index
binding_lengths variadic of index

hal.device.queue.execute (HAL::DeviceQueueExecuteOp)link

Enqueues command buffer execution

Syntax:

operation ::= `hal.device.queue.execute` `<` $device `:` type($device) `>`
              `affinity` `(` $queue_affinity `)`
              `wait` `(` $wait_fence `)`
              `signal` `(` $signal_fence `)`
              (`commands` `(` `[` $command_buffers^ `]` `)`)?
              attr-dict-with-keyword

Executes one or more command buffers on a device queue. The command buffers are executed in order as if they were recorded as one. No commands will execute until the wait fence has been reached and the signal fence will be signaled when all commands have completed.

Operands:link
Operand Description
device device
queue_affinity 64-bit signless integer
wait_fence fence
signal_fence fence
command_buffers variadic of command_buffer

hal.device.queue.flush (HAL::DeviceQueueFlushOp)link

Flushes locally-pending submissions to the queue

Syntax:

operation ::= `hal.device.queue.flush` `<` $device `:` type($device) `>`
              `affinity` `(` $queue_affinity `)`
              attr-dict-with-keyword

Flushes any locally-pending submissions in the queue. When submitting many queue operations this can be used to eagerly flush earlier submissions while later ones are still being constructed. This may be a no-op.

Operands:link
Operand Description
device device
queue_affinity 64-bit signless integer

hal.device.queue.read (HAL::DeviceQueueReadOp)link

Reads a segment from a file into a device buffer

Syntax:

operation ::= `hal.device.queue.read` `<` $device `:` type($device) `>`
              `affinity` `(` $queue_affinity `)`
              `wait` `(` $wait_fence `)`
              `signal` `(` $signal_fence `)`
              `source` `(` $source_file `:` type($source_file) `)`
              `` `[` $source_offset `]`
              `target` `(` $target_buffer `:` type($target_buffer) `)`
              `` `[` $target_offset `]`
              `length` `(` $length `)`
              `flags` `(` $flags `)`
              attr-dict-with-keyword

Enqueues a file read operation that streams a segment of the source file defined by the source offset and length into the target HAL buffer at the specified target offset. The queue affinity should be set to where the target buffer will be consumed. The source file must have read permission and the target buffer must have transfer-target usage. Read failure will result in propagated semaphore failure or device loss.

Attributes:link
AttributeMLIR TypeDescription
flags::mlir::IntegerAttr32-bit signless integer attribute
Operands:link
Operand Description
device device
queue_affinity 64-bit signless integer
wait_fence fence
signal_fence fence
source_file buffer
source_offset 64-bit signless integer
target_buffer buffer
target_offset index
length index

hal.device.queue.write (HAL::DeviceQueueWriteOp)link

Writes a segment from a device buffer into a file

Syntax:

operation ::= `hal.device.queue.write` `<` $device `:` type($device) `>`
              `affinity` `(` $queue_affinity `)`
              `wait` `(` $wait_fence `)`
              `signal` `(` $signal_fence `)`
              `source` `(` $source_buffer `:` type($source_buffer) `)`
              `` `[` $source_offset `]`
              `target` `(` $target_file `:` type($target_file) `)`
              `` `[` $target_offset `]`
              `length` `(` $length `)`
              `flags` `(` $flags `)`
              attr-dict-with-keyword

Enqueues a file write operation that streams a segment of the source HAL buffer defined by the source offset and length into the target file at the specified target offset. The queue affinity should be set to where the source buffer was produced. The source buffer must have transfer-source usage and the target file must have write permission. Write failure will result in propagated semaphore failure or device loss.

Attributes:link
AttributeMLIR TypeDescription
flags::mlir::IntegerAttr32-bit signless integer attribute
Operands:link
Operand Description
device device
queue_affinity 64-bit signless integer
wait_fence fence
signal_fence fence
source_buffer buffer
source_offset index
target_file buffer
target_offset 64-bit signless integer
length index

hal.return (HAL::ReturnOp)link

Return from a hal.* region

Syntax:

operation ::= `hal.return` ($operands^ `:` type($operands))? attr-dict

Returns the given values from the region and back to the host code.

Traits: Terminator

Operands:link
Operand Description
operands variadic of any type

Executable opslink

Ops for !hal.executable / iree_hal_executable_t.

hal.executable.binary (HAL::ExecutableBinaryOp)link

Compiled executable binary data

Syntax:

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

A compiled executable binary with an optional nested module containing the IR prior to serialization (for debugging).

Traits: HasParent<IREE::HAL::ExecutableOp>

Interfaces: Symbol

Attributes:link
AttributeMLIR TypeDescription
sym_visibility::mlir::StringAttrstring attribute
sym_name::mlir::StringAttrstring attribute
format::mlir::StringAttrstring attribute
data::mlir::DenseIntElementsAttr8-bit signless integer elements attribute
mime_type::mlir::StringAttrstring attribute

hal.executable.calculate_workgroups (HAL::ExecutableCalculateWorkgroupsOp)link

Calculates workgroup count from workload for an exported function

Syntax:

operation ::= `hal.executable.calculate_workgroups` `device` `(` $device `:` type($device) `)`
              `target` `(` $entry_point `)`
              (`workload` `(` `[` $workload^ `]` `)`)?
              `:` type($workgroup_x) `,` type($workgroup_y) `,` type($workgroup_z)
              attr-dict-with-keyword

Calculates the workgroup count (grid XYZ) based on the given workload using the workgroup count calculation region of the target hal.executable.export op.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:link
AttributeMLIR TypeDescription
entry_point::mlir::SymbolRefAttrsymbol reference attribute
Operands:link
Operand Description
device device
workload variadic of index
Results:link
Result Description
workgroup_x index
workgroup_y index
workgroup_z index

hal.executable.condition (HAL::ExecutableConditionOp)link

Host code to determine if the executable is enabled

Variants are selected based on their target and this optional condition op that returns true if the variant is valid for use on the provided runtime !hal.device. If no variants within an executable are valid then loading will fail at runtime. If multiple variants are valid the first valid one found will be loaded and used for execution.

Traits: IsolatedFromAbove

Interfaces: CallableOpInterface, FunctionOpInterface, Symbol

Attributes:link
AttributeMLIR TypeDescription
function_type::mlir::TypeAttrtype attribute of function type
arg_attrs::mlir::ArrayAttrArray of dictionary attributes
res_attrs::mlir::ArrayAttrArray of dictionary attributes

hal.executable.constant.block (HAL::ExecutableConstantBlockOp)link

Executable constant block initializer

Initializes one or more constants in the executable constant block by returning one value per identified constant. Each constant block is evaluated on the host prior to instantiating the executable for a given device and allows for the executable to be specialized based on device capabilities and limits.

The keys specified are unique per variant and will be deduplicated across multiple constant blocks when present. They are only used during lowering and will not survive to runtime so they need only have descriptive enough names to avoid collisions and represent the semantics of the value.

Constant values can be loaded in the device code with the hal.executable.constant.load op:

hal.executable.variant public @target {
  hal.executable.constant.block(%device: !hal.device) -> (i32, i32) as ("foo", "bar") {
    %0 = hal.device.query<%device> key("some.device.prop")...
    %1 = hal.device.query<%device> key("another.device.prop")...
    hal.return %0, %1 : i32, i32
  }
  builtin.module {
    func @dispatch0() {
      %0 = hal.executable.constant.load "foo" : i32
      %1 = hal.executable.constant.load "bar" : i32
      return
    }
  }
}

Each target backend will implement the constant initialization and access in a way compatible with its execution model. Examples: - CPU: read-only buffer initialized on load and passed to each dispatch - CUDA: read-only buffer initialized on load and passed to each dispatch - SPIR-V: specialization constants - Metal: function constants - WebGPU: pipeline-overridable constants

Traits: HasParent<IREE::HAL::ExecutableSourceOp, IREE::HAL::ExecutableVariantOp>, IsolatedFromAbove

Interfaces: CallableOpInterface, FunctionOpInterface, Symbol

Attributes:link
AttributeMLIR TypeDescription
function_type::mlir::TypeAttrtype attribute of function type
keys::mlir::ArrayAttrarray attribute
arg_attrs::mlir::ArrayAttrArray of dictionary attributes
res_attrs::mlir::ArrayAttrArray of dictionary attributes

hal.executable.constant.load (HAL::ExecutableConstantLoadOp)link

Loads a constant value from the executable constant block

Syntax:

operation ::= `hal.executable.constant.load` $key attr-dict `:` type($result)

Loads a scalar constant value from the static executable constant block. The value provided by a constant block with the given key will be loaded and bitcast (possibly with truncation or zero-extension) to the result type.

Note that backends are allowed to implement their own mechanisms for referencing constant block values and this is provided only as a default for those not needing special behavior.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:link
AttributeMLIR TypeDescription
key::mlir::StringAttrstring attribute
Results:link
Result Description
result index or signless integer or floating-point or complex-type

hal.executable.create (HAL::ExecutableCreateOp)link

Creates an executable

Syntax:

operation ::= `hal.executable.create` `device` `(` $device `:` type($device) `)`
              `target` `(` $executable_target `)`
              `layouts` `(` `[` $layouts `]` `)`
              (`constants` `(` `[` $constants^ `]` `)`)?
              `:` type($result)
              attr-dict-with-keyword

Creates a target-dependent executable cached on the provided device. Entry points contained within the executable can be dispatched using the resulting executable handle.

Depending on the driver creation may take a non-trivial amount of time (such as when JITing/etc). As the cache is internally synchronized callers can issue preparation requests from multiple threads - even for the same executables - and calls will block until preparation completes.

Optional constants provide for specialization of the executable based on runtime-derived parameters.

Traits: AlwaysSpeculatableImplTrait, AttrSizedOperandSegments

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface), OpAsmOpInterface

Effects: MemoryEffects::Effect{}

Attributes:link
AttributeMLIR TypeDescription
executable_target::mlir::SymbolRefAttrsymbol reference attribute
Operands:link
Operand Description
device device
layouts variadic of pipeline_layout
constants variadic of 32-bit signless integer
Results:link
Result Description
result executable

hal.executable_end (HAL::ExecutableEndOp)link

Terminator pseudo-op for the executable op

Syntax:

operation ::= `hal.executable_end` attr-dict

Traits: HasParent<IREE::HAL::ExecutableOp>, Terminator

hal.executable.export (HAL::ExecutableExportOp)link

Executable entry point declaration

An entry point exported by the executable with statically-available information describing the IO interface it uses and other dispatch metadata.

The workgroup_count region represents the computation that returns the number of workgroups to use in the 3D grid dispatch. The arguments to the region represents the workload as captured by each dispatch. It returns the number of workgroups along x, y, and z.

Traits: HasParent<IREE::HAL::ExecutableSourceOp, IREE::HAL::ExecutableVariantOp>, 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::HAL::PipelineLayoutAttrexecutable entry point layout specification
workgroup_size::mlir::ArrayAttrindex array attribute
subgroup_size::mlir::IntegerAttrsize_t
workgroup_local_memory::mlir::IntegerAttrindex attribute
source_locs::mlir::DictionaryAttrdictionary of named attribute values

hal.executable.export.ordinal (HAL::ExecutableExportOrdinalOp)link

Executable export ordinal lookup pseudo-op

Syntax:

operation ::= `hal.executable.export.ordinal` `target` `(` $entry_point `)`
              `:` type($result)
              attr-dict-with-keyword

Resolves an executable export ordinal to a value once ordinals have been assigned.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface), OpAsmOpInterface

Effects: MemoryEffects::Effect{}

Attributes:link
AttributeMLIR TypeDescription
entry_point::mlir::SymbolRefAttrsymbol reference attribute
Results:link
Result Description
result index

hal.executable.lookup (HAL::ExecutableLookupOp)link

Executable cache lookup pseudo-op

Syntax:

operation ::= `hal.executable.lookup` `device` `(` $device `:` type($device) `)`
              `executable` `(` $executable `)`
              `:` type($result)
              attr-dict-with-keyword

Used during conversion to provide a placeholder for a globally cached and possibly lazy-initialized executable.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface), OpAsmOpInterface

Effects: MemoryEffects::Effect{}

Attributes:link
AttributeMLIR TypeDescription
executable::mlir::FlatSymbolRefAttrflat symbol reference attribute
Operands:link
Operand Description
device device
Results:link
Result Description
result executable

hal.executable (HAL::ExecutableOp)link

Target-specific executable module

Syntax:

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

An executable module representing a target-specific compiled kernel/shader/etc.

Traits: IsolatedFromAbove, SingleBlockImplicitTerminator<IREE::HAL::ExecutableEndOp>, SingleBlock, SymbolTable, Util_ObjectLike

Interfaces: Symbol

Attributes:link
AttributeMLIR TypeDescription
sym_visibility::mlir::StringAttrstring attribute
sym_name::mlir::StringAttrstring attribute

hal.executable.source_end (HAL::ExecutableSourceEndOp)link

Terminator pseudo-op for the executable source op

Syntax:

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

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

hal.executable.source (HAL::ExecutableSourceOp)link

Generic source contents of an executable op

Syntax:

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

This is an unspecialized source representation of an executable module without an assigned target. This is useful for hand-authoring executables prior to device specification.

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

Interfaces: Symbol

Attributes:link
AttributeMLIR TypeDescription
sym_visibility::mlir::StringAttrstring attribute
sym_name::mlir::StringAttrstring attribute
objects::mlir::iree_compiler::IREE::HAL::ExecutableObjectsAttrtarget-specific object file references

hal.executable.variant_end (HAL::ExecutableVariantEndOp)link

Terminator pseudo-op for the executable variant op

Syntax:

operation ::= `hal.executable.variant_end` attr-dict

Traits: HasParent<IREE::HAL::ExecutableVariantOp>, Terminator

hal.executable.variant (HAL::ExecutableVariantOp)link

Target-specific variant of an executable op

Syntax:

operation ::= `hal.executable.variant` custom<SymbolVisibility>($sym_visibility)
              $sym_name
              `target` `(` $target `)`
              (`objects` `(` $objects^ `)` )?
              (`sources` `(` $sources^ `)` )?
              attr-dict-with-keyword
              $body

The target IR for the executable. This can be preserved for debugging but is usually removed during transformation.

Variants are selected based on their target and an optional condition op that returns true if the variant is valid for use on the provided runtime !hal.device. If no variants within an executable are valid then loading will fail at runtime. If multiple variants are valid the first valid one found will be loaded and used for execution.

Traits: HasParent<IREE::HAL::ExecutableOp>, IsolatedFromAbove, SingleBlockImplicitTerminator<IREE::HAL::ExecutableVariantEndOp>, SingleBlock, SymbolTable

Interfaces: Symbol

Attributes:link
AttributeMLIR TypeDescription
sym_visibility::mlir::StringAttrstring attribute
sym_name::mlir::StringAttrstring attribute
target::mlir::iree_compiler::IREE::HAL::ExecutableTargetAttrgeneric executable target specification
objects::mlir::ArrayAttrHAL executable object references
sources::mlir::DictionaryAttrdictionary of named attribute values

Experimental opslink

Temporary hack ops expected to be removed in the future.

hal.ex.file.from_memory (HAL::ExFileFromMemoryOp)link

Creates a file mapped into a byte range of a host buffer

Syntax:

operation ::= `hal.ex.file.from_memory` `device` `(` $device `:` type($device) `)`
              `affinity` `(` $queue_affinity `)`
              `access` `(` $access `)`
              `buffer` `(` $buffer `:` type($buffer) `)`
              `` `[` $offset `for` $length `]`
              `flags` `(` $flags `)`
              `:` type($result)
              attr-dict-with-keyword

Returns a file handle that is backed by the given buffer contents. Behavior is undefined if the buffer contents change while the accesses are in-flight.

Experimental as the exact interface for getting files from module contents still needs iteration. Most hardware APIs require a file descriptor or native platform handle but here we only have host pointers. When memory-mapped some systems allow for retrieval of the platform handle from a virtual address (GetMappedFileNameA/posix_mem_offset) but the APIs are sketchy and likely slow. Instead we should probably have a way to query for a file handle derived from the calling module by stack-walking and asking the VM module for its handle. Until we can figure this out this method will be marked epxerimental.

Interfaces: OpAsmOpInterface

Attributes:link
AttributeMLIR TypeDescription
accessmlir::iree_compiler::IREE::HAL::MemoryAccessBitfieldAttrvalid MemoryAccess
Operands:link
Operand Description
device device
queue_affinity 64-bit signless integer
buffer a reference counted byte buffer
offset index
length index
flags 32-bit signless integer
Results:link
Result Description
result buffer

Fence opslink

Ops for !hal.fence / iree_hal_fence_t.

hal.fence.await (HAL::FenceAwaitOp)link

Asynchronous fence wait operation

Syntax:

operation ::= `hal.fence.await` `until` `(` `[` $fences `]` `)`
              `timeout_millis` `(` $timeout_millis `)`
              `:` type($status)
              attr-dict-with-keyword

Yields the caller until all fences is reached. Returns the status of the fence after the wait, with a non-zero value indicating failure.

Traits: Util_YieldPoint

Interfaces: OpAsmOpInterface

Operands:link
Operand Description
timeout_millis 32-bit signless integer
fences variadic of fence
Results:link
Result Description
status 32-bit signless integer

hal.fence.create (HAL::FenceCreateOp)link

Creates an unsignaled fence

Syntax:

operation ::= `hal.fence.create` `device` `(` $device `:` type($device) `)`
              `flags` `(` $flags `)`
              `:` type($result)
              attr-dict-with-keyword

Returns a fence that defines a point in time. By default fences will remain unsignaled unless they are explicitly signaled with hal.fence.signal or asynchronously signaled by the device by passing them as an operand to queue submission ops.

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface), OpAsmOpInterface

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

Attributes:link
AttributeMLIR TypeDescription
flagsmlir::iree_compiler::IREE::HAL::FenceFlagBitfieldAttrvalid FenceFlag
Operands:link
Operand Description
device device
Results:link
Result Description
result fence

hal.fence.fail (HAL::FenceFailOp)link

Fence failure operation

Syntax:

operation ::= `hal.fence.fail` `<` $fence `:` type($fence) `>`
              `status` `(` $status `)`
              attr-dict-with-keyword

Signals the fence with a failure. The status will be returned from each timepoint semaphores hal.semaphore.query and hal.semaphore.signal for the lifetime of each semaphore.

Operands:link
Operand Description
fence fence
status 32-bit signless integer

hal.fence.join (HAL::FenceJoinOp)link

Creates a fence from the given timepoints

Syntax:

operation ::= `hal.fence.join` `at` `(` `[` $fences `]` `)`
              `->` type($result)
              attr-dict-with-keyword

Returns a fence that joins the input fences as a wait-all operation.

Interfaces: OpAsmOpInterface

Operands:link
Operand Description
fences variadic of fence
Results:link
Result Description
result fence

hal.fence.query (HAL::FenceQueryOp)link

Fence query operation

Syntax:

operation ::= `hal.fence.query` `<` $fence `:` type($fence) `>`
              `:` type($status)
              attr-dict-with-keyword

Queries whether the fence has been reached and its status. Returns OK if the fence has been signaled successfully, DEFERRED if it is unsignaled, and otherwise an error indicating the failure.

Operands:link
Operand Description
fence fence
Results:link
Result Description
status 32-bit signless integer

hal.fence.signal (HAL::FenceSignalOp)link

Fence signal operation

Syntax:

operation ::= `hal.fence.signal` `<` $fence `:` type($fence) `>`
              attr-dict-with-keyword

Signals the fence to indicate that the timepoints contained have been reached. Waiting work may begin immediately.

Operands:link
Operand Description
fence fence

Instrument opslink

Ops for !hal.instrument.*.

hal.instrument.memory.load (HAL::InstrumentMemoryLoadOp)link

Emits a memory load instrumentation event

Syntax:

operation ::= `hal.instrument.memory.load` `` `[` $buffer `:` type($buffer) `for` $workgroupKey `]`
              $base `[` $indices `]` `,` $loadValue
              attr-dict `:` type($base) `,` type($result)

Emits a workgroup-specific memory load event indicating that a number of bytes from the given resolved pointer have been loaded by the workgroup.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:link
Operand Description
buffer memref of any type values
workgroupKey index
loadValue any type
base memref of any type values
indices variadic of index
Results:link
Result Description
result any type

hal.instrument.memory.store (HAL::InstrumentMemoryStoreOp)link

Emits a memory store instrumentation event

Syntax:

operation ::= `hal.instrument.memory.store` `` `[` $buffer `:` type($buffer) `for` $workgroupKey `]`
              $base `[` $indices `]` `,` $storeValue
              attr-dict `:` type($base) `,` type($result)

Emits a workgroup-specific memory store event indicating that a number of bytes have been stored to the given resolved pointer by the workgroup.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:link
Operand Description
buffer memref of any type values
workgroupKey index
storeValue any type
base memref of any type values
indices variadic of index
Results:link
Result Description
result any type

hal.instrument.print (HAL::InstrumentPrintOp)link

Emits a human-readable printf-style string event

Syntax:

operation ::= `hal.instrument.print` `` `[` $buffer `:` type($buffer) `for` $workgroupKey `]`
              $format (`*` `(` $values^ `:` type($values) `)`)?
              attr-dict

Formats a string using a limited subset of printf format specifiers and the provided values and then emits an iree_instrument_dispatch_print_t event. Final formatted string lengths may be limited to as much as 1024 characters and should be kept as small as possible to avoid easily exceeding the instrumentation storage buffers with redundant strings.

Attributes:link
AttributeMLIR TypeDescription
format::mlir::StringAttrstring attribute
Operands:link
Operand Description
buffer memref of any type values
workgroupKey index
values variadic of any type

hal.instrument.value (HAL::InstrumentValueOp)link

Emits a scalar value instrumentation event

Syntax:

operation ::= `hal.instrument.value` `` `[` $buffer `:` type($buffer) `for` $workgroupKey `]`
              $ordinal `=` $operand attr-dict `:` type($operand)

Emits a workgroup-specific typed value with the given workgroup-relative ordinal.

This op will be preserved even if the output is not used as it is only for debugging purposes.

Attributes:link
AttributeMLIR TypeDescription
ordinal::mlir::IntegerAttr8-bit integer attribute
Operands:link
Operand Description
buffer memref of any type values
workgroupKey index
operand any type
Results:link
Result Description
result any type

hal.instrument.workgroup (HAL::InstrumentWorkgroupOp)link

Emits a dispatch workgroup instrumentation event

Syntax:

operation ::= `hal.instrument.workgroup` `` `[` $buffer `:` type($buffer) `]`
              `dispatch` `(` $dispatchId `)`
              attr-dict `:` type($workgroupKey)

Emits an iree_instrument_dispatch_workgroup_t event into the instrumentation stream. The workgroup event identifies the unique dispatch, its workgroup count, and the ID of the emitting workgroup within the dispatch. Optionally targets that support querying the processor ID executing the workgroup can attach that information for tracking purposes.

On targets such as CPUs where entire workgroups execute as atomic units only one workgroup event should be emitted. On targets such as GPUs where there may be multiple invocations executing as part of a single workgroup only the first invocation within the workgroup should emit the workgroup event (by checking if the LocalInvocationIndex or threadIdx == 0, etc).

The resulting workgroup key is used by subsequent workgroup-specific instrumentation events.

Operands:link
Operand Description
buffer memref of any type values
dispatchId 32-bit signless integer
Results:link
Result Description
workgroupKey index

Interface opslink

Ops for !hal.interface.*.

hal.interface.binding.subspan (HAL::InterfaceBindingSubspanOp)link

Returns an alias to a subspan of interface binding data

Syntax:

operation ::= `hal.interface.binding.subspan` `set` `(` $set `)`
              `binding` `(` $binding `)`
              `type` `(` custom<DescriptorType>($descriptor_type) `)`
              (`alignment` `(` $alignment^ `)`)?
              (`offset` `(` $byte_offset^ `)`)?
              (`flags` `(` $descriptor_flags^ `)`)?
              attr-dict `:` type($result) (`{` $dynamic_dims^ `}`)?

Returns a subspan of an interface binding storage buffer in a generic type. The exact shape, type, and alignment of the returned type are defined by the result type (tensor, memref, etc).

An optional alignment indicates the byte alignment of the base binding resource. Note that the byte offset is added to the base and the alignment will be the minimum of the two.

Traits: AlwaysSpeculatableImplTrait, AttrSizedOperandSegments

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface), Util_ShapeAwareOp

Effects: MemoryEffects::Effect{}

Attributes:link
AttributeMLIR TypeDescription
set::mlir::IntegerAttrindex attribute
binding::mlir::IntegerAttrindex attribute
descriptor_type::mlir::iree_compiler::IREE::HAL::DescriptorTypeAttrvalid DescriptorType
alignment::mlir::IntegerAttrindex attribute
descriptor_flags::mlir::iree_compiler::IREE::HAL::DescriptorFlagsAttrvalid Descriptor flags
Operands:link
Operand Description
byte_offset index
dynamic_dims variadic of index
Results:link
Result Description
result any type

hal.interface.constant.load (HAL::InterfaceConstantLoadOp)link

Loads a constant value from the interface constant block

Syntax:

operation ::= `hal.interface.constant.load` `` `[` $index `]`
              (`alignment` `(` $alignment^ `)`)?
              (`values` `(` $values^ `)`)?
              attr-dict `:` type($result)

Loads a scalar constant value from an executable IO push constant block. The value will be loaded from the given constant offset and will be bitcast (possibly with truncation or zero-extension) to the result type.

An optional alignment indicates the byte alignment of potential values for the constant when it could be determined from analysis. If omitted the value may be anything and its interpretation is up to the usage. This is intended to provide pointer alignment-like semantics to constants that are used to index into binding resources.

An optional set of values indicates all possible values that can be passed to the constant from all dispatch sites in the program. If omitted the value may be from an unanalyzable source (outside of the program, indirect, etc) and must be assumed to have any value.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:link
AttributeMLIR TypeDescription
index::mlir::IntegerAttrsize_t
alignment::mlir::IntegerAttrindex attribute
values::mlir::ArrayAttrarray attribute
Results:link
Result Description
result index or signless integer or floating-point or complex-type

hal.interface.workgroup.count (HAL::InterfaceWorkgroupCountOp)link

Returns the total workgroup count of the grid

Syntax:

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

The total number of workgroups along each dimension in the dispatch grid. Matches what was passed to the hal.command_buffer.dispatch command (or what was indirectly specified).

Corresponds to the NumWorkgroups SPIR-V built-in and the gridDim CUDA built-in variable.

%x = hal.interface.workgroup.count[0] : index
%y = hal.interface.workgroup.count[1] : index
%z = hal.interface.workgroup.count[2] : index

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface), OpAsmOpInterface

Effects: MemoryEffects::Effect{}

Attributes:link
AttributeMLIR TypeDescription
dimension::mlir::IntegerAttrindex attribute
Results:link
Result Description
result index

hal.interface.workgroup.id (HAL::InterfaceWorkgroupIDOp)link

Returns the index of the current workgroup in the grid

Syntax:

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

The global workgroup ID of the current tile in the range of [0, hal.interface.workgroup.count) along each XYZ dimension.

Corresponds to the WorkgroupId SPIR-V built-in and the blockIdx CUDA built-in variable.

%x = hal.interface.workgroup.id[0] : index
%y = hal.interface.workgroup.id[1] : index
%z = hal.interface.workgroup.id[2] : index

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface), OpAsmOpInterface

Effects: MemoryEffects::Effect{}

Attributes:link
AttributeMLIR TypeDescription
dimension::mlir::IntegerAttrindex attribute
Results:link
Result Description
result index

hal.interface.workgroup.size (HAL::InterfaceWorkgroupSizeOp)link

Returns the size of each workgroup in invocations

Syntax:

operation ::= `hal.interface.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.

Corresponds to the WorkgroupSize SPIR-V built-in and the blockDim CUDA built-in variable.

%x = hal.interface.workgroup.size[0] : index
%y = hal.interface.workgroup.size[1] : index
%z = hal.interface.workgroup.size[2] : index

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface), OpAsmOpInterface

Effects: MemoryEffects::Effect{}

Attributes:link
AttributeMLIR TypeDescription
dimension::mlir::IntegerAttrindex attribute
Results:link
Result Description
result index

Pipeline layout opslink

Ops for !hal.pipeline_layout / iree_hal_pipeline_layout_t.

hal.pipeline_layout.create (HAL::PipelineLayoutCreateOp)link

Creates an pipeline layout

Syntax:

operation ::= `hal.pipeline_layout.create` `device` `(` $device `:` type($device) `)`
              `push_constants` `(` $push_constants `)`
              `layouts` `(` `[` $set_layouts `]` `)`
              `:` type($result)
              attr-dict-with-keyword

Creates an pipeline layout from the given descriptor sets and push constant required size. Pipeline layouts can be shared across any executable that uses the same layout and push constant information. Sharing the layout between executables will reduce runtime binding overhead and it is often worth the cost to allow a small number of unused bindings in one executable such that it can share layouts with others that will be scheduled adjacent to it.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface), OpAsmOpInterface

Effects: MemoryEffects::Effect{}

Attributes:link
AttributeMLIR TypeDescription
push_constants::mlir::IntegerAttrindex attribute
Operands:link
Operand Description
device device
set_layouts variadic of descriptor_set_layout
Results:link
Result Description
result pipeline_layout

hal.pipeline_layout.lookup (HAL::PipelineLayoutLookupOp)link

Pipeline layout cache lookup pseudo-op

Syntax:

operation ::= `hal.pipeline_layout.lookup` `device` `(` $device `:` type($device) `)`
              `layout` `(` $layout `)`
              `:` type($result)
              attr-dict-with-keyword

Used during conversion to provide a placeholder for a globally cached and possibly lazy-initialized pipeline layout.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface), OpAsmOpInterface

Effects: MemoryEffects::Effect{}

Attributes:link
AttributeMLIR TypeDescription
layout::mlir::iree_compiler::IREE::HAL::PipelineLayoutAttrexecutable entry point layout specification
Operands:link
Operand Description
device device
Results:link
Result Description
result pipeline_layout

Pseudo Opslink

Pseudo ops for conversion support.

hal.dispatch.extern (HAL::DispatchExternOp)link

A dispatch of workgroups across a 3-dimensional grid

Syntax:

operation ::= `hal.dispatch.extern` $export
              (`[` $workload^ `]`)? ``
              `(` $arguments `)` `:`
              custom<ShapedFunctionType>(ref($arguments),
              type($arguments), $argument_dims,
              type($results), $result_dims,
              $tied_operands)
              `count` `` custom<WorkgroupCountRegion>($workgroup_count)
              `layout` `(` $layout `)`
              (`bindings` `(` $bindings^ `)`)?
              `objects` `(` `{` custom<TargetConditionObjects>($targets,
              $target_ordinals,
              $target_objects,
              $target_regions) `}` `)`
              attr-dict-with-keyword

Dispatches some number of workgroups across a 3-dimensional grid using a function defined externally in one or more referenced objects. Objects are declared per executable target and selected automatically during linking based on where the dispatch is used. Semantically this is equivalent to a flow.dispatch.workgroups but with the workgroup region invisible to the compiler. See hal.executable for more information about object linkage.

Note that since this happens at tensor level the dispatch operation has value semantics: some tensors (and optionally other primitive types) are consumed and one or more new result tensors are produced. Inside each workgroup, however, the input and output tensors are available for arbitrary loads and stores. In many cases each workgroup will load some particular tile(s) from the input tensors and store some particular tile(s) to the output tensors unique to that workgroup. Though it's possible for multiple workgroups to load the same regions of the input tensors behavior is undefined if multiple workgroups store to the same regions of the output tensors. Codegen guarantees this behavior but when sourcing externally authored dispatch functions it's critical that this behavior is observed.

Though the representation is similar to the GPU-style grid dispatch model here we still have not yet allocated buffers, determined the target device for execution, or even completed fully resolving shapes/types/etc. Because of this it's important that the workgroup body use the platform-dependent primitives for accessing workgroup ID, size, and count intrinsics instead of hardcoding them to a particular set of values. Assume that any workgroup dispatch may end up being specialized for several different target devices and even several different variants for a particular target device (differing workgroup sizes, etc). To aid deduplication code producing these external dispatches should try not to specialize early for particular shapes and instead emit the most generic code possible as having 500 slightly different hal.dispatch.extern ops pointing at the same object file is likely to require 500 copies of the object instead of 500 calls to the same object.

Because at this point in the layering devices have not yet been selected the workgroup count cannot be fully evaluated. Instead workload parameters are captured that are then passed to a function that when later evaluated computes the actual workgroup count based on target information. The workload is not limited to the 3D XYZ grid dispatch of the workgroup count and can contain any number of parameters used to compute it. If workgroup size or distribution varies based on the target device a !hal.device argument can be used by the workgroup count calculation region to factor in device parameters. See hal.device.query for more information on how to query information.

%r = hal.dispatch.extern "some_function"[%c5, %c5](%0, %1)
    : (tensor<5x5xf32>, tensor<5xf32>) -> tensor<5x5xf32>
  ...

The number of results of the operation is equal to the number of results in the type signature ((tensor<5x5xf32>, tensor<5xf32>) -> tensor<5x5xf32>). Each tensor argument and result in the type signature has a corresponding pipeline layout slot and must be declared. If multiple arguments or results share the same layout slot they can be aliased using the bindings attribute and otherwise each is assumed unique.

There are no arguments operands for results, but a result can be tied an argument by writing the argument operand's SSA value instead of its type: E.g., in the above example, -> %0 would tie the first argument to the result. In that case, there would be no separate block argument for the result.

Objects for multiple targets can be specified and the ones used are selected based on their target and an optional condition region that returns true if the variant is valid for use on the provided runtime !hal.device. If no variants within an executable are valid then loading will fail at runtime. If multiple variants are valid the first valid one found will be loaded and used for execution.

Traits: AlwaysSpeculatableImplTrait, AttrSizedOperandSegments, IsolatedFromAbove

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

Effects: MemoryEffects::Effect{}

Attributes:link
AttributeMLIR TypeDescription
export::mlir::StringAttrstring attribute
layout::mlir::iree_compiler::IREE::HAL::PipelineLayoutAttrexecutable entry point layout specification
targets::mlir::ArrayAttrarray attribute
target_ordinals::mlir::ArrayAttrArray of index ordinal attributes
target_objects::mlir::ArrayAttrarray attribute
workgroup_size::mlir::ArrayAttrindex array attribute
subgroup_size::mlir::IntegerAttrsize_t
workgroup_local_memory::mlir::IntegerAttrindex attribute
bindings::mlir::ArrayAttrHAL binding array 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

hal.tensor.alias (HAL::TensorAliasOp)link

Hints that tensor storage should alias a HAL buffer view

Syntax:

operation ::= `hal.tensor.alias` (`wait` `(` $wait_fence^ `)` `=` `` `>`)?
              $source `:` type($source) (`{` $source_dims^ `}`)?
              `to`
              $storage `:` type($storage)
              attr-dict

Hints that the backing storage of an entire tensor aliases the given storage buffer. There's no guarantee that the storage will alias and instead only that the tensor contents will be written to the storage as if a copy had occurred. This allows the compiler to avoid copies in the ideal case of a producer that is able to produce directly into the target storage but still handle cases where the producer is not able to be in-place.

The storage buffer provided must have sufficient space for the tensor once encoded. Dynamically shaped tensors may not consume the entire provided storage. If a buffer view is provided the metadata is ignored and only the backing buffer is used.

An optional wait fence can be provided in cases where the storage is not immediately available. Producers that may alias the storage will wait until the storage is available before updating the contents.

Explicit aliasing side-steps any analysis that may be performed by the compiler and requires users to guarantee that the safety of the aliasing. Copy-on-write, alias analysis for overlap detection, and ordering via use-def chains are all ignorant of the aliased buffer memory and only ensure the compiler consumes or produces the aliased memory consistent with itself.

Example:

%init = tensor.empty
%value = linalg.generic ... outs(%init)
%aliased = hal.tensor.alias %value : tensor<...> to %buffer : !hal.buffer
... linalg.generic ins(%aliased) ...

Traits: AlwaysSpeculatableImplTrait, AttrSizedOperandSegments

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

Effects: MemoryEffects::Effect{}

Operands:link
Operand Description
source tensor of any type values
source_dims variadic of index
storage buffer or buffer_view
wait_fence fence
Results:link
Result Description
result tensor of any type values

hal.tensor.barrier (HAL::TensorBarrierOp)link

Signals a fence when all tensors are available

Syntax:

operation ::= `hal.tensor.barrier` `join` `` `(` $sources `:` type($sources) `)`
              `=` `` `>`
              $signal_fence `:` type($signal_fence)
              attr-dict-with-keyword

Defines a barrier that is used to indicate availability of an entire set of tensors by signaling a fence. The source tensors are returned for chaining.

Interfaces: TiedOpInterface

Operands:link
Operand Description
sources variadic of tensor of any type values
signal_fence fence
Results:link
Result Description
results variadic of tensor of any type values

hal.tensor.export (HAL::TensorExportOp)link

Exports a tensor to a HAL buffer view

Syntax:

operation ::= `hal.tensor.export` $source
              ($name^)?
              `:`
              custom<TypeAlias>($source_encoding, type($source)) (`{` $source_dims^ `}`)?
              `->`
              type($target)
              attr-dict

Defines an export of an SSA-form tensor to an external HAL buffer view.

The provided source_encoding, if different from the source type, indicates that the ABI-facing type may differ from the internal representation. The types must be bitcastable (same storage size) and dynamically shaped values must have the same number of dynamic dimensions. This allows for casting between rank-0 and rank-N types, different element types, etc.

Traits: AlwaysSpeculatableImplTrait

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

Effects: MemoryEffects::Effect{}

Attributes:link
AttributeMLIR TypeDescription
source_encoding::mlir::TypeAttrany type attribute
name::mlir::StringAttrstring attribute
Operands:link
Operand Description
source tensor of any type values
source_dims variadic of index
Results:link
Result Description
target buffer or buffer_view

hal.tensor.import (HAL::TensorImportOp)link

Imports a tensor from a HAL buffer view

Syntax:

operation ::= `hal.tensor.import` (`wait` `(` $wait_fence^ `)` `=` `` `>`)?
              $source
              ($name^)?
              `:` type($source) `->`
              custom<TypeAlias>($target_encoding, type($target)) (`{` $target_dims^ `}`)?
              attr-dict

Defines an import of an external HAL buffer view into a SSA-form tensor. An optional semaphore timepoint can be specified indicating when the buffer view is available for use. If no semaphore timepoint is provided it is assumed the buffer view is immediately available.

The provided target_encoding, if different from the target type, indicates that the ABI-facing type may differ from the internal representation. The types must be bitcastable (same storage size) and dynamically shaped values must have the same number of dynamic dimensions. This allows for casting between rank-0 and rank-N types, different element types, etc.

Traits: AlwaysSpeculatableImplTrait, AttrSizedOperandSegments

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

Effects: MemoryEffects::Effect{}

Attributes:link
AttributeMLIR TypeDescription
target_encoding::mlir::TypeAttrany type attribute
name::mlir::StringAttrstring attribute
Operands:link
Operand Description
source buffer or buffer_view
target_dims variadic of index
wait_fence fence
Results:link
Result Description
target tensor of any type values

Attributeslink

AffinityQueueAttrlink

specifies a set of allowed queues for an operation

WIP; see #10765. This may change in the future to either be a nested attribute on a larger affinity struct or be defined by an implementation of the affinity attr interface. For now this allows higher levels of the stack to specify queues such that the stream dialect can understand them and they can be lowered into the HAL dialect.

Specifies that an annotated operation or scope is only allowed to execute on the set of queues (0-64) provided. Operations will not run on other queues.

Example:

// any queue
#hal.affinity.queue<*>
// queues 4 and 5
#hal.affinity.queue<[4, 5]>

Parameters:link
Parameter C++ type Description
mask int64_t

CollectiveAttrlink

collective operation and specification

Syntax:

#hal.collective<
  CollectiveKind,   # kind
  std::optional<CollectiveReductionOp>,   # reduction
  CollectiveElementType   # element_type
>

Specifies the collective operation to perform and any mode bits required.

Parameters:link
Parameter C++ type Description
kind CollectiveKind
reduction std::optional<CollectiveReductionOp>
element_type CollectiveElementType

DescriptorSetBindingAttrlink

descriptor set binding specification

Syntax:

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

Specifies a single binding within a descriptor set layout.

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

DescriptorSetLayoutAttrlink

descriptor set layout specification

Syntax:

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

Specifies the layout information of a single set of descriptors used within an pipeline layout. Multiple of these sets may be used by a single entry point to allow for bindings with similar update frequencies to be grouped.

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

DescriptorTypeAttrlink

valid DescriptorType

Syntax:

#hal.descriptor_type<
  ::mlir::iree_compiler::IREE::HAL::DescriptorType   # value
>

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

Parameters:link
Parameter C++ type Description
value ::mlir::iree_compiler::IREE::HAL::DescriptorType an enum of type DescriptorType

DeviceTargetAttrlink

generic device target specification

Specifies the properties of a target runtime device. Target devices are specified with a canonical identifier matching those used by the runtime (such as cpu, vulkan, etc). Target devices may support several target executable formats specified with #hal.executable.target. An optional configuration dictionary allows for overriding backend defaults.

Example:

#hal.device.target<"llvm-cpu", {
  device_configuration = ...
}, [
  #hal.executable.target<"llvm-cpu", "embedded-elf-arm_32">,
  #hal.executable.target<"llvm-cpu", "embedded-elf-arm_64">,
]>

Parameters:link
Parameter C++ type Description
deviceID StringAttr
configuration DictionaryAttr
executable_targets ::llvm::ArrayRef<ExecutableTargetAttr>

ExecutableObjectAttrlink

object file reference

Defines an object file that can be linked into executables. Today this is only supported for external file references with paths the compiler can successfully resolve from its current working directory. Inlined data can optionally be provided to avoid the need for file system access and ensure the data source is attached to the IR as it makes its way through multiple compiler stages or reproducers.

Future revisions may change this to an interface that allows both internal and external resources to define the object contents. Linking needs to be updated to support various object compositions and certain backends may require additional infrastructure support.

In the long term the goal is to allow combinations of declared objects and generated code in order to give control of linking behavior to frontends. Instead of needing global command line flags to link in additional blobs the frontend can emit executables with the dependencies already defined per variant without needing to reach into the IREE compiler code.

Example:

#hal.executable.object<{path = "some/file.obj"}>
#hal.executable.object<{
  path = "some/embedded/file.obj",
  data = dense<[...]> : vector<2048xi8>
}>

Parameters:link
Parameter C++ type Description
path StringAttr
data DenseIntElementsAttr

ExecutableObjectsAttrlink

target-specific object file references

A dictionary mapping executable target specifications to a list of objects. This is used to allow layers of the stack that support multi-targeting to specify information used during lowering into each particular target.

The key attributes are matched against each target variant based on the backend and format as well as any configuration data provided. When comparing the configuration only fields present in both the key and target variant will be checked and must match. This allows specification of generic sets ("all x86_64 targets get these objects") as well as specific ones ("only x86_64 targets with vector_size = 64 get these objects").

Example:

#hal.executable.objects<{
  #hal.executable.target<"llvm-cpu", "embedded-elf-arm_64"> = [
    #hal.executable.object<{path = "some/file_arm_64.obj"}>
  ],
  #hal.executable.target<"llvm-cpu", "embedded-elf-x86_64"> = [
    #hal.executable.object<{path = "some/file_x86_64.obj"}>
  ]
}>

Parameters:link
Parameter C++ type Description
targets ArrayAttr
targetObjects ArrayAttr

ExecutableTargetAttrlink

generic executable target specification

Specifies how to compile an executable for a specific target backend. A backend is used to translate and serialize the executable into the final form passed to the runtime. The format of the executable is a target-specific value indicating the required runtime support to load the deployed artifact. An optionally provided configuration dictionary overrides backend-specific defaults.

Example:

  // Produce a system-native ELF for x86-64 systems using the LLVM backend:
  #hal.executable.target<"llvm-cpu", "system-elf-x86_64", {
    triple = "x86_64-unknown-linux-elf",
    cpu = "host",
    cpu_features = "host",
    abi = "lp32",
    ...
  }>

The same compilation backend may be used to translate executables for several different runtime devices. Likewise the same runtime device may use one of many different executable targets. Assume an N:M mapping between the two in all cases.

Parameters:link
Parameter C++ type Description
backend StringAttr
format StringAttr
configuration DictionaryAttr

InterfaceBindingAttrlink

interface binding specification

Syntax:

#hal.interface.binding<
  int64_t,   # set
  int64_t   # binding
>

Specifies the descriptor set and binding ordinal of a particular layout binding.

Example:

#hal.interface.binding<0, 1>

Parameters:link
Parameter C++ type Description
set int64_t
binding int64_t

PipelineLayoutAttrlink

executable entry point layout specification

Syntax:

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

Specifies the layout information used for interacting with executable functions. This allows host code to correctly map parameters to the lower-level target-specific argument passing behavior.

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

Type constraintslink

allocatorlink

Allocates buffers for a particular device memory space.

bufferlink

A memory buffer with a specific memory_type that is used to describe the capabilities and behavior of the backing memory of the buffer. Buffers may be any mix of host-accessible, host-coherent, or device-accessible for various usages. Depending on these memory types the buffers may be mapped for access on the host as memory though certain restrictions may be imposed.

buffer_viewlink

A shaped and typed buffer reference. This just wraps an existing hal.buffer with its associated metadata to make it easier to pass across ABI boundaries. In most cases buffer views can be elided entirely by the compiler and they'll only be seen when calling external functions.

collective.channellink

Channel identifier used to allow for participation in multiple collective groups.

command_bufferlink

Asynchronous command buffer recording interface. Commands are recorded by the implementation for later submission to command queues.

descriptor_set_layoutlink

Descriptor set layout.

devicelink

Logical device instance.

eventlink

Events are used for defining synchronization scopes within CommandBuffers. An event only exists within a single CommandBuffer and must not be used across CommandBuffers from the same device or others.

executablelink

A prepared and ready-to-dispatch executable.

fencelink

A set of semaphore timepoints defining a common point in time across multiple timelines.

bufferlink

A stateless file handle that can be read/written using queue-ordered transfer operations.

pipeline_layoutlink

A pipeline layout describing the descriptor sets and push constants used.