TritonInstrumentOps¶

`tti.experimental_assert_in_thread` (triton::instrument::ExperimentalAssertInThreadOp)¶

Assert the condition within the current thread

Syntax:

operation ::= `tti.experimental_assert_in_thread` $condition `,` $message attr-dict `:` type($condition)

Assert that the condition is true given all the values are available in the current thread. If the condition is false, the message is printed, and the program is aborted. If check_any is true, any of the values in the condition must be true. Otherwise, all the values in the condition must be true.

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}

Attributes:¶

Attribute	MLIR Type	Description
`message`	::mlir::StringAttr	string attribute
`check_any`	::mlir::BoolAttr	bool attribute

Operands:¶

Operand	Description
`condition`	tensor of 1-bit signless integer values

`tti.experimental_check_outstanding_reads` (triton::instrument::ExperimentalCheckOutstandingReadsOp)¶

Check if there are outstanding reads from a buffer guarded by a mbar

Syntax:

operation ::= `tti.experimental_check_outstanding_reads` $buf `{` $buffers `,` $readBars `(` $readBarsType `)` `}` (`,` $pred^)? attr-dict `:` type($buf) `,` type($buffers) `,` type($readBars)

Check if there are outstanding reads from a buffer guarded by a mbar.

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}

Attributes:¶

Attribute	MLIR Type	Description
`readBarsType`	::mlir::TypeAttr	any type attribute

Operands:¶

Operand	Description
`buf`	memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system
`buffers`	ranked tensor of floating-point or integer or ptr values
`readBars`	ptr or ranked tensor of ptr values
`pred`	1-bit signless integer

`tti.experimental_check_outstanding_writes` (triton::instrument::ExperimentalCheckOutstandingWritesOp)¶

Check if there are outstanding writes to a buffer guarded by a mbar

Syntax:

operation ::= `tti.experimental_check_outstanding_writes` $buf `{` $buffers `,` $writeBars `(` $writeBarsType `)` `}` (`,` $pred^)? attr-dict `:` type($buf) `,` type($buffers) `,` type($writeBars)

Check if there are outstanding writes to a buffer guarded by a mbar.

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}

Attributes:¶

Attribute	MLIR Type	Description
`writeBarsType`	::mlir::TypeAttr	any type attribute

Operands:¶

Operand	Description
`buf`	memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system
`buffers`	ranked tensor of floating-point or integer or ptr values
`writeBars`	ptr or ranked tensor of ptr values
`pred`	1-bit signless integer

`tti.experimental_check_write_commit` (triton::instrument::ExperimentalCheckWriteCommitOp)¶

Check if the buffer has an outstanding write commit.

Syntax:

operation ::= `tti.experimental_check_write_commit` $buf `{` $buffers `,` $writeCommits `(` $writeCommitsType `)` `}` (`,` $pred^)? attr-dict `:` type($buf) `,` type($buffers) `,` type($writeCommits)

Check if the buffer has an outstanding write commit.

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}

Attributes:¶

Attribute	MLIR Type	Description
`writeCommitsType`	::mlir::TypeAttr	any type attribute

Operands:¶

Operand	Description
`buf`	memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system
`buffers`	ranked tensor of floating-point or integer or ptr values
`writeCommits`	ptr or ranked tensor of ptr values
`pred`	1-bit signless integer

`tti.experimental_clear_read_barrier` (triton::instrument::ExperimentalClearReadBarrierOp)¶

Clear the read state for buffers being guarded by an mbar

Syntax:

operation ::= `tti.experimental_clear_read_barrier` $mbar `{` $barriers `,` $readBars `(` $readBarsType `)` `}` (`,` $pred^)? attr-dict `:` type($mbar) `,` type($barriers) `,` type($readBars)

Clear the read state for buffers being guarded by an mbar.

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}

Attributes:¶

Attribute	MLIR Type	Description
`readBarsType`	::mlir::TypeAttr	any type attribute

Operands:¶

Operand	Description
`mbar`	memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system
`barriers`	ranked tensor of floating-point or integer or ptr values
`readBars`	ptr or ranked tensor of ptr values
`pred`	1-bit signless integer

`tti.experimental_clear_write_barrier` (triton::instrument::ExperimentalClearWriteBarrierOp)¶

Clear the write state for buffers being guarded by an mbar

Syntax:

operation ::= `tti.experimental_clear_write_barrier` $mbar `{` $writeBars `(` $writeBarsType `)` `}` (`,` $pred^)? attr-dict `:` type($mbar) `,` type($writeBars)

Clear the write state for buffers being guarded by an mbar.

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}

Attributes:¶

Attribute	MLIR Type	Description
`writeBarsType`	::mlir::TypeAttr	any type attribute

Operands:¶

Operand	Description
`mbar`	memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system
`writeBars`	ptr or ranked tensor of ptr values
`pred`	1-bit signless integer

`tti.experimental_clear_write_commits` (triton::instrument::ExperimentalClearWriteCommitsOp)¶

Clear all the write commits more distant than `outstandingNum.

Syntax:

operation ::= `tti.experimental_clear_write_commits` `{` $writeCommits `(` $writeCommitsType `)` `}` `,` $outstandingNum (`,` $pred^)? attr-dict `:` type($writeCommits)

Clear all the write commits more distant than outstandingNum from the current thread.

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}

Attributes:¶

Attribute	MLIR Type	Description
`writeCommitsType`	::mlir::TypeAttr	any type attribute
`outstandingNum`	::mlir::IntegerAttr	32-bit signless integer attribute

Operands:¶

Operand	Description
`writeCommits`	ptr or ranked tensor of ptr values
`pred`	1-bit signless integer

`tti.experimental_commit_writes` (triton::instrument::ExperimentalCommitWritesOp)¶

Commit all the staged writes for all the buffers.

Syntax:

operation ::= `tti.experimental_commit_writes` `{` $writeCommits `(` $writeCommitsType `)` `}` (`,` $pred^)? attr-dict `:` type($writeCommits)

Commit all the staged writes for all the buffers.

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}

Attributes:¶

Attribute	MLIR Type	Description
`writeCommitsType`	::mlir::TypeAttr	any type attribute

Operands:¶

Operand	Description
`writeCommits`	ptr or ranked tensor of ptr values
`pred`	1-bit signless integer

`tti.experimental_mark_as_read` (triton::instrument::ExperimentalMarkAsReadOp)¶

Mark a buffer as being read from using mbar as a guard

Syntax:

operation ::= `tti.experimental_mark_as_read` $buf `,` $mbar `{` $buffers `,` $barriers `,` $readBars `(` $readBarsType `)` `}` (`,` $pred^)? attr-dict `:` type($buf) `,` type($mbar) `,` type($buffers) `,` type($barriers) `,` type($readBars)

Mark a buffer as being read from using mbar as a guard.

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}

Attributes:¶

Attribute	MLIR Type	Description
`readBarsType`	::mlir::TypeAttr	any type attribute

Operands:¶

Operand	Description
`buf`	memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system
`mbar`	memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system
`buffers`	ranked tensor of floating-point or integer or ptr values
`barriers`	ranked tensor of floating-point or integer or ptr values
`readBars`	ptr or ranked tensor of ptr values
`pred`	1-bit signless integer

`tti.experimental_mark_as_write` (triton::instrument::ExperimentalMarkAsWriteOp)¶

Mark a buffer as being written to using mbar as a guard

Syntax:

operation ::= `tti.experimental_mark_as_write` $buf `,` $mbar `{` $buffers `,` $writeBars `(` $writeBarsType `)` `}` (`,` $pred^)? attr-dict `:` type($buf) `,` type($mbar) `,` type($buffers) `,` type($writeBars)

Mark a buffer as being written to using mbar as a guard.

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}

Attributes:¶

Attribute	MLIR Type	Description
`writeBarsType`	::mlir::TypeAttr	any type attribute

Operands:¶

Operand	Description
`buf`	memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system
`mbar`	memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system
`buffers`	ranked tensor of floating-point or integer or ptr values
`writeBars`	ptr or ranked tensor of ptr values
`pred`	1-bit signless integer

`tti.experimental_shared_buffer_pointers` (triton::instrument::ExperimentalSharedBufferPointersOp)¶

Definte an array of pointers to shared memory buffers

Syntax:

operation ::= `tti.experimental_shared_buffer_pointers` attr-dict `:` type($result)

Create a tensor of pointers to shared memory buffers.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:¶

Attribute	MLIR Type	Description
`offsets`	::mlir::DenseI32ArrayAttr	i32 dense array attribute

Results:¶

Result	Description
`result`	ranked tensor of floating-point or integer or ptr values

`tti.experimental_stage_write_for_commit` (triton::instrument::ExperimentalStageWriteForCommitOp)¶

Preapre to an async copy of a buffer. Staged until commit_group is called.

Syntax:

operation ::= `tti.experimental_stage_write_for_commit` $buf `{` $buffers `,` $writeCommits `(` $writeCommitsType `)` `}` (`,` $pred^)? attr-dict `:` type($buf) `,` type($buffers) `,` type($writeCommits)

Preapre to an async copy of a buffer. Staged until commit_group is called. The implementation will write -1 to the write_commits tensor under the indices corresponding to the buffer.

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}

Attributes:¶

Attribute	MLIR Type	Description
`writeCommitsType`	::mlir::TypeAttr	any type attribute

Operands:¶

Operand	Description
`buf`	memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system
`buffers`	ranked tensor of floating-point or integer or ptr values
`writeCommits`	ptr or ranked tensor of ptr values
`pred`	1-bit signless integer

TritonInstrumentOps¶

tti.experimental_assert_in_thread (triton::instrument::ExperimentalAssertInThreadOp)¶

Attributes:¶

Operands:¶

tti.experimental_check_outstanding_reads (triton::instrument::ExperimentalCheckOutstandingReadsOp)¶

Attributes:¶

Operands:¶

tti.experimental_check_outstanding_writes (triton::instrument::ExperimentalCheckOutstandingWritesOp)¶

Attributes:¶

Operands:¶

tti.experimental_check_write_commit (triton::instrument::ExperimentalCheckWriteCommitOp)¶

Attributes:¶

Operands:¶

tti.experimental_clear_read_barrier (triton::instrument::ExperimentalClearReadBarrierOp)¶

Attributes:¶

Operands:¶

tti.experimental_clear_write_barrier (triton::instrument::ExperimentalClearWriteBarrierOp)¶

Attributes:¶

Operands:¶

tti.experimental_clear_write_commits (triton::instrument::ExperimentalClearWriteCommitsOp)¶

Attributes:¶

Operands:¶

tti.experimental_commit_writes (triton::instrument::ExperimentalCommitWritesOp)¶

Attributes:¶

Operands:¶

tti.experimental_mark_as_read (triton::instrument::ExperimentalMarkAsReadOp)¶

Attributes:¶

Operands:¶

tti.experimental_mark_as_write (triton::instrument::ExperimentalMarkAsWriteOp)¶

Attributes:¶

Operands:¶

tti.experimental_shared_buffer_pointers (triton::instrument::ExperimentalSharedBufferPointersOp)¶

Attributes:¶

Results:¶

tti.experimental_stage_write_for_commit (triton::instrument::ExperimentalStageWriteForCommitOp)¶

Attributes:¶

Operands:¶

`tti.experimental_assert_in_thread` (triton::instrument::ExperimentalAssertInThreadOp)¶

`tti.experimental_check_outstanding_reads` (triton::instrument::ExperimentalCheckOutstandingReadsOp)¶

`tti.experimental_check_outstanding_writes` (triton::instrument::ExperimentalCheckOutstandingWritesOp)¶

`tti.experimental_check_write_commit` (triton::instrument::ExperimentalCheckWriteCommitOp)¶

`tti.experimental_clear_read_barrier` (triton::instrument::ExperimentalClearReadBarrierOp)¶

`tti.experimental_clear_write_barrier` (triton::instrument::ExperimentalClearWriteBarrierOp)¶

`tti.experimental_clear_write_commits` (triton::instrument::ExperimentalClearWriteCommitsOp)¶

`tti.experimental_commit_writes` (triton::instrument::ExperimentalCommitWritesOp)¶

`tti.experimental_mark_as_read` (triton::instrument::ExperimentalMarkAsReadOp)¶

`tti.experimental_mark_as_write` (triton::instrument::ExperimentalMarkAsWriteOp)¶

`tti.experimental_shared_buffer_pointers` (triton::instrument::ExperimentalSharedBufferPointersOp)¶

`tti.experimental_stage_write_for_commit` (triton::instrument::ExperimentalStageWriteForCommitOp)¶