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 |
|---|---|
|
tensor of 1-bit signless integer values |
tti.experimental_buffer_pointers (triton::instrument::ExperimentalBufferPointersOp)¶
Definte an array of pointers to shared memory buffers
Syntax:
operation ::= `tti.experimental_buffer_pointers` $offsets `,` $memType 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 |
memType | ::mlir::triton::instrument::MemTypeAttr | allowed 32-bit signless integer cases: 0, 1 |
Results:¶
Result |
Description |
|---|---|
|
ranked tensor of floating-point or integer or ptr values |
tti.experimental_check_barrier_writes_cleared (triton::instrument::ExperimentalCheckBarrierWritesClearedOp)¶
Verify that the barrier is not used to track any writes
Syntax:
operation ::= `tti.experimental_check_barrier_writes_cleared` $mbar `{` $barriers `,` $writeBars `(` $writeBarsType `)` `}` (`,` $pred^)? attr-dict `:` type($mbar) `,` type($barriers) `,` type($writeBars)
Verify that the barrier is not used to track any writes.
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 |
|---|---|
|
memory descriptor type ( |
|
ranked tensor of floating-point or integer or ptr values |
|
ptr or ranked tensor of ptr values |
|
1-bit signless integer |
tti.experimental_check_outstanding_commits (triton::instrument::ExperimentalCheckOutstandingCommitsOp)¶
Check if the buffer has an outstanding commit.
Syntax:
operation ::= `tti.experimental_check_outstanding_commits` $buf `{` $buffers `,` $outstandingCommits `(` $outstandingCommitsType `)` `}` `,` $pendingAccessType (`,` $pred^)? attr-dict `:` type($buf) `,` type($buffers) `,` type($outstandingCommits)
Check if the buffer has an outstanding commit.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Attributes:¶
| Attribute | MLIR Type | Description |
|---|---|---|
outstandingCommitsType | ::mlir::TypeAttr | any type attribute |
pendingAccessType | ::mlir::StringAttr | string attribute |
Operands:¶
Operand |
Description |
|---|---|
|
memory descriptor type ( |
|
ranked tensor of floating-point or integer or ptr values |
|
ptr or ranked tensor of ptr values |
|
1-bit signless integer |
tti.experimental_check_read_barriers (triton::instrument::ExperimentalCheckReadBarriersOp)¶
Check if there are outstanding reads from a buffer guarded by a mbar
Syntax:
operation ::= `tti.experimental_check_read_barriers` $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 |
|---|---|
|
memory descriptor type ( |
|
ranked tensor of floating-point or integer or ptr values |
|
ptr or ranked tensor of ptr values |
|
1-bit signless integer |
tti.experimental_check_write_state (triton::instrument::ExperimentalCheckWriteStateOp)¶
Check if there are outstanding writes to a buffer guarded by a mbar
Syntax:
operation ::= `tti.experimental_check_write_state` $buf `{` $buffers `,` $writeBars `(` $writeBarsType `)` `,` $writeState `(` $writeStateType `)` `}` (`,` $pred^)? `pipelined` $hwPipelined attr-dict `:` type($buf) `,` type($buffers) `,` type($writeBars) `,` type($writeState)
Check if the writeState tensor has non-zero value associated with the buffer.
writeState is a tensor of 8b bitfields, where:
bit 0: 1 if the buffer is being written to
bit 1: 1 if the write is not hwPipelined
If hwPipelined is true, shift the bitfield by 1 to check the second bit - this means that the error won’t be triggered if another pipelined write is outstanding.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Attributes:¶
| Attribute | MLIR Type | Description |
|---|---|---|
writeBarsType | ::mlir::TypeAttr | any type attribute |
writeStateType | ::mlir::TypeAttr | any type attribute |
hwPipelined | ::mlir::IntegerAttr | 1-bit signless integer attribute |
Operands:¶
Operand |
Description |
|---|---|
|
memory descriptor type ( |
|
ranked tensor of floating-point or integer or ptr values |
|
ptr or ranked tensor of ptr values |
|
ptr or ranked tensor of ptr values |
|
1-bit signless integer |
tti.experimental_clear_outstanding_commits (triton::instrument::ExperimentalClearOutstandingCommitsOp)¶
Clear all the outstanding commits more distant than `outstandingNum.
Syntax:
operation ::= `tti.experimental_clear_outstanding_commits` `{` $outstandingCommits `(` $outstandingCommitsType `)` `}` `,` $outstandingNum (`,` $pred^)? attr-dict `:` type($outstandingCommits)
Clear all the outstanding 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 |
|---|---|---|
outstandingCommitsType | ::mlir::TypeAttr | any type attribute |
outstandingNum | ::mlir::IntegerAttr | 32-bit signless integer attribute |
Operands:¶
Operand |
Description |
|---|---|
|
ptr or ranked tensor of ptr values |
|
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 |
|---|---|
|
memory descriptor type ( |
|
ranked tensor of floating-point or integer or ptr values |
|
ptr or ranked tensor of ptr values |
|
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 `{` $barriers `,` $writeBars `(` $writeBarsType `)` `,` $writeState `(` $writeStateType `)` `}` (`,` $pred^)? attr-dict `:` type($mbar) `,` type($barriers) `,` type($writeBars) `,` type($writeState)
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 |
writeStateType | ::mlir::TypeAttr | any type attribute |
Operands:¶
Operand |
Description |
|---|---|
|
memory descriptor type ( |
|
ranked tensor of floating-point or integer or ptr values |
|
ptr or ranked tensor of ptr values |
|
ptr or ranked tensor of ptr values |
|
1-bit signless integer |
tti.experimental_commit_accesses (triton::instrument::ExperimentalCommitAccessesOp)¶
Commit all the staged accesses for all the buffers.
Syntax:
operation ::= `tti.experimental_commit_accesses` `{` $outstandingCommits `(` $outstandingCommitsType `)` `}` (`,` $pred^)? attr-dict `:` type($outstandingCommits)
Commit all the staged accesses for all the buffers.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Attributes:¶
| Attribute | MLIR Type | Description |
|---|---|---|
outstandingCommitsType | ::mlir::TypeAttr | any type attribute |
Operands:¶
Operand |
Description |
|---|---|
|
ptr or ranked tensor of ptr values |
|
1-bit signless integer |
tti.experimental_commit_write_with_barrier (triton::instrument::ExperimentalCommitWriteWithBarrierOp)¶
Mark all buffers being currently written as tracked by the barrier.
Syntax:
operation ::= `tti.experimental_commit_write_with_barrier` $mbar `{` $barriers `,` $writeBars `(` $writeBarsType `)` `,` $writeState `(` $writeStateType `)` `}` (`,` $pred^)? attr-dict `:` type($mbar) `,` type($barriers) `,` type($writeBars) `,` type($writeState)
For all buffers currently marked in writeState tensor, mark them as tracked by the mbar in writeBars tensor.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Attributes:¶
| Attribute | MLIR Type | Description |
|---|---|---|
writeBarsType | ::mlir::TypeAttr | any type attribute |
writeStateType | ::mlir::TypeAttr | any type attribute |
Operands:¶
Operand |
Description |
|---|---|
|
memory descriptor type ( |
|
ranked tensor of floating-point or integer or ptr values |
|
ptr or ranked tensor of ptr values |
|
ptr or ranked tensor of ptr values |
|
1-bit signless integer |
tti.experimental_set_read_barrier (triton::instrument::ExperimentalSetReadBarrierOp)¶
Mark a buffer as being read from using mbar as a guard
Syntax:
operation ::= `tti.experimental_set_read_barrier` $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 |
|---|---|
|
memory descriptor type ( |
|
memory descriptor type ( |
|
ranked tensor of floating-point or integer or ptr values |
|
ranked tensor of floating-point or integer or ptr values |
|
ptr or ranked tensor of ptr values |
|
1-bit signless integer |
tti.experimental_set_write_state (triton::instrument::ExperimentalSetWriteStateOp)¶
Mark a buffer as being written in writeState tensor
Syntax:
operation ::= `tti.experimental_set_write_state` $buf `{` $buffers `,` $writeState `(` $writeStateType `)` `}` (`,` $pred^)? `pipelined` $hwPipelined attr-dict `:` type($buf) `,` type($buffers) `,` type($writeState)
Mark a buffer as being written to. It is not yet tracked by a barrier, until
commit_write_with_barrier is called, at which point all the buffers being written
to are marked as tracked by the barrier.
writeState is a tensor of 8b bitfields, where:
bit 0: 1 if the buffer is being written to
bit 1: 1 if the write is not hwPipelined
If hwPipelined is true, the write won’t trigger an error if another pipelined write is executed later without waiting for the barrier.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Attributes:¶
| Attribute | MLIR Type | Description |
|---|---|---|
writeStateType | ::mlir::TypeAttr | any type attribute |
hwPipelined | ::mlir::IntegerAttr | 1-bit signless integer attribute |
Operands:¶
Operand |
Description |
|---|---|
|
memory descriptor type ( |
|
ranked tensor of floating-point or integer or ptr values |
|
ptr or ranked tensor of ptr values |
|
1-bit signless integer |
tti.experimental_stage_access_for_commit (triton::instrument::ExperimentalStageAccessForCommitOp)¶
Syntax:
operation ::= `tti.experimental_stage_access_for_commit` $buf `{` $buffers `,` $outstandingCommits `(` $outstandingCommitsType `)` `}` (`,` $pred^)? attr-dict `:` type($buf) `,` type($buffers) `,` type($outstandingCommits)
For operations that use outstanding to track the number of outstanding commits (rather than mbarriers),
mark the buffer as being accessed, but not commited yet, by marking it with -1.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Attributes:¶
| Attribute | MLIR Type | Description |
|---|---|---|
outstandingCommitsType | ::mlir::TypeAttr | any type attribute |
Operands:¶
Operand |
Description |
|---|---|
|
memory descriptor type ( |
|
ranked tensor of floating-point or integer or ptr values |
|
ptr or ranked tensor of ptr values |
|
1-bit signless integer |