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

'hal' Dialect

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

Attribute	MLIR Type	Description
`memory_types`	mlir::iree_compiler::IREE::HAL::MemoryTypeBitfieldAttr	valid MemoryType
`buffer_usage`	mlir::iree_compiler::IREE::HAL::BufferUsageBitfieldAttr	valid 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

Attribute	MLIR Type	Description
`memory_types`	mlir::iree_compiler::IREE::HAL::MemoryTypeBitfieldAttr	valid MemoryType
`buffer_usage`	mlir::iree_compiler::IREE::HAL::BufferUsageBitfieldAttr	valid 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

Attribute	MLIR Type	Description
`message`	::mlir::StringAttr	string attribute
`memory_types`	mlir::iree_compiler::IREE::HAL::MemoryTypeBitfieldAttr	valid MemoryType
`buffer_usage`	mlir::iree_compiler::IREE::HAL::BufferUsageBitfieldAttr	valid 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

Attribute	MLIR Type	Description
`message`	::mlir::StringAttr	string 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

Attribute	MLIR Type	Description
`index`	::mlir::IntegerAttr	index 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

Attribute	MLIR Type	Description
`key`	::mlir::StringAttr	string 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

Attribute	MLIR Type	Description
`type`	::mlir::TypeAttr	any 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

Attribute	MLIR Type	Description
`encoding`	::mlir::Attribute	any 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

Attribute	MLIR Type	Description
`flags`	::mlir::IntegerAttr	32-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

Attribute	MLIR Type	Description
`flags`	::mlir::IntegerAttr	32-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

Attribute	MLIR Type	Description
`label`	::mlir::StringAttr	string 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

Attribute	MLIR Type	Description
`op`	::mlir::iree_compiler::IREE::HAL::CollectiveAttr	collective 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

Attribute	MLIR Type	Description
`modes`	mlir::iree_compiler::IREE::HAL::CommandBufferModeBitfieldAttr	valid CommandBufferMode
`command_categories`	mlir::iree_compiler::IREE::HAL::CommandCategoryBitfieldAttr	valid 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

Attribute	MLIR Type	Description
`source_stage_mask`	mlir::iree_compiler::IREE::HAL::ExecutionStageBitfieldAttr	valid ExecutionStage
`target_stage_mask`	mlir::iree_compiler::IREE::HAL::ExecutionStageBitfieldAttr	valid ExecutionStage
`flags`	mlir::iree_compiler::IREE::HAL::ExecutionBarrierFlagBitfieldAttr	valid 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

Attribute	MLIR Type	Description
`offset`	::mlir::IntegerAttr	index 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

Attribute	MLIR Type	Description
`flags`	::mlir::iree_compiler::IREE::HAL::DescriptorSetLayoutFlagsAttr	valid DescriptorSetLayout flags
`bindings`	::mlir::ArrayAttr	HAL 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

Attribute	MLIR Type	Description
`category`	::mlir::StringAttr	string attribute
`key`	::mlir::StringAttr	string attribute
`default_value`	::mlir::TypedAttr	TypedAttr 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

Attribute	MLIR Type	Description
`memory_types`	mlir::iree_compiler::IREE::HAL::MemoryTypeBitfieldAttr	valid MemoryType
`buffer_usage`	mlir::iree_compiler::IREE::HAL::BufferUsageBitfieldAttr	valid 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` (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

Attribute	MLIR Type	Description
`flags`	::mlir::IntegerAttr	32-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

Attribute	MLIR Type	Description
`flags`	::mlir::IntegerAttr	32-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

Attribute	MLIR Type	Description
`sym_visibility`	::mlir::StringAttr	string attribute
`sym_name`	::mlir::StringAttr	string attribute
`format`	::mlir::StringAttr	string attribute
`data`	::mlir::DenseIntElementsAttr	8-bit signless integer elements attribute
`mime_type`	::mlir::StringAttr	string 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

Attribute	MLIR Type	Description
`entry_point`	::mlir::SymbolRefAttr	symbol 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

Attribute	MLIR Type	Description
`function_type`	::mlir::TypeAttr	type attribute of function type
`arg_attrs`	::mlir::ArrayAttr	Array of dictionary attributes
`res_attrs`	::mlir::ArrayAttr	Array 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

Attribute	MLIR Type	Description
`function_type`	::mlir::TypeAttr	type attribute of function type
`keys`	::mlir::ArrayAttr	array attribute
`arg_attrs`	::mlir::ArrayAttr	Array of dictionary attributes
`res_attrs`	::mlir::ArrayAttr	Array 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

Attribute	MLIR Type	Description
`key`	::mlir::StringAttr	string 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

Attribute	MLIR Type	Description
`executable_target`	::mlir::SymbolRefAttr	symbol 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

Attribute	MLIR Type	Description
`sym_visibility`	::mlir::StringAttr	string attribute
`sym_name`	::mlir::StringAttr	string attribute
`ordinal`	::mlir::IntegerAttr	size_t
`layout`	::mlir::iree_compiler::IREE::HAL::PipelineLayoutAttr	executable entry point layout specification
`workgroup_size`	::mlir::ArrayAttr	index array attribute
`subgroup_size`	::mlir::IntegerAttr	size_t
`workgroup_local_memory`	::mlir::IntegerAttr	index attribute
`source_locs`	::mlir::DictionaryAttr	dictionary 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

Attribute	MLIR Type	Description
`entry_point`	::mlir::SymbolRefAttr	symbol 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

Attribute	MLIR Type	Description
`executable`	::mlir::FlatSymbolRefAttr	flat 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

Attribute	MLIR Type	Description
`sym_visibility`	::mlir::StringAttr	string attribute
`sym_name`	::mlir::StringAttr	string 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

Attribute	MLIR Type	Description
`sym_visibility`	::mlir::StringAttr	string attribute
`sym_name`	::mlir::StringAttr	string attribute
`objects`	::mlir::iree_compiler::IREE::HAL::ExecutableObjectsAttr	target-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

Attribute	MLIR Type	Description
`sym_visibility`	::mlir::StringAttr	string attribute
`sym_name`	::mlir::StringAttr	string attribute
`target`	::mlir::iree_compiler::IREE::HAL::ExecutableTargetAttr	generic executable target specification
`objects`	::mlir::ArrayAttr	HAL executable object references
`sources`	::mlir::DictionaryAttr	dictionary 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

Attribute	MLIR Type	Description
`access`	mlir::iree_compiler::IREE::HAL::MemoryAccessBitfieldAttr	valid 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

Attribute	MLIR Type	Description
`flags`	mlir::iree_compiler::IREE::HAL::FenceFlagBitfieldAttr	valid 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

Attribute	MLIR Type	Description
`format`	::mlir::StringAttr	string 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

Attribute	MLIR Type	Description
`ordinal`	::mlir::IntegerAttr	8-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

Attribute	MLIR Type	Description
`set`	::mlir::IntegerAttr	index attribute
`binding`	::mlir::IntegerAttr	index attribute
`descriptor_type`	::mlir::iree_compiler::IREE::HAL::DescriptorTypeAttr	valid DescriptorType
`alignment`	::mlir::IntegerAttr	index attribute
`descriptor_flags`	::mlir::iree_compiler::IREE::HAL::DescriptorFlagsAttr	valid 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

Attribute	MLIR Type	Description
`index`	::mlir::IntegerAttr	size_t
`alignment`	::mlir::IntegerAttr	index attribute
`values`	::mlir::ArrayAttr	array 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

Attribute	MLIR Type	Description
`dimension`	::mlir::IntegerAttr	index 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

Attribute	MLIR Type	Description
`dimension`	::mlir::IntegerAttr	index 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

Attribute	MLIR Type	Description
`dimension`	::mlir::IntegerAttr	index 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

Attribute	MLIR Type	Description
`push_constants`	::mlir::IntegerAttr	index 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

Attribute	MLIR Type	Description
`layout`	::mlir::iree_compiler::IREE::HAL::PipelineLayoutAttr	executable 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

Attribute	MLIR Type	Description
`export`	::mlir::StringAttr	string attribute
`layout`	::mlir::iree_compiler::IREE::HAL::PipelineLayoutAttr	executable entry point layout specification
`targets`	::mlir::ArrayAttr	array attribute
`target_ordinals`	::mlir::ArrayAttr	Array of index ordinal attributes
`target_objects`	::mlir::ArrayAttr	array attribute
`workgroup_size`	::mlir::ArrayAttr	index array attribute
`subgroup_size`	::mlir::IntegerAttr	size_t
`workgroup_local_memory`	::mlir::IntegerAttr	index attribute
`bindings`	::mlir::ArrayAttr	HAL binding array attribute
`tied_operands`	::mlir::ArrayAttr	64-bit integer array attribute

Operands:link

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

Results:link

Result	Description
`results`	variadic of any type

`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

Attribute	MLIR Type	Description
`source_encoding`	::mlir::TypeAttr	any type attribute
`name`	::mlir::StringAttr	string 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

Attribute	MLIR Type	Description
`target_encoding`	::mlir::TypeAttr	any type attribute
`name`	::mlir::StringAttr	string 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.

'hal' Dialectlink

Operationslink

Allocator opslink

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

Attributes:link

Operands:link

Results:link

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

Attributes:link

Operands:link

Results:link

Buffer opslink

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

Attributes:link

Operands:link

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

Operands:link

Results:link

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

Operands:link

Results:link

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

Operands:link

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

Operands:link

Results:link

Buffer view opslink

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

Attributes:link

Operands:link

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

Operands:link

Results:link

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

Operands:link

Results:link

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

Attributes:link

Operands:link

Results:link

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

Operands:link

Results:link

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

Operands:link

Results:link

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

Operands:link

Results:link

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

Attributes:link

Operands:link

hal.element_type (HAL::ElementTypeOp)link

Attributes:link

Results:link

hal.encoding_type (HAL::EncodingTypeOp)link

Attributes:link

Results:link

Channel opslink

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

Attributes:link

Operands:link

Results:link

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

Operands:link

Results:link

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

Attributes:link

Operands:link

Results:link

Command buffer opslink

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

Attributes:link

Operands:link

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

Attributes:link

Operands:link

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

Operands:link

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

`hal.element_type` (HAL::ElementTypeOp)link

`hal.encoding_type` (HAL::EncodingTypeOp)link

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

`hal.return` (HAL::ReturnOp)link

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

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

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