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_all_active_waiting (triton::instrument::ExperimentalCheckAllActiveWaitingOp)¶
Assert that not all active threads are waiting on matching phases
Syntax:
operation ::= `tti.experimental_check_all_active_waiting` $activeMask `,` $barriers `,` $waiting `(` $waitingType `)` `,` $barrierStates `(` $barrierStatesType `)` (`,` $pred^)? attr-dict `:` type($barriers) `,` type($waiting) `,` type($barrierStates)
Filter waiting threads to those whose recorded phase matches the current barrier phase, OR-reduce across barriers, and assert that (waitingMask & activeMask) != activeMask.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Attributes:¶
| Attribute | MLIR Type | Description |
|---|---|---|
activeMask | ::mlir::IntegerAttr | 32-bit signless integer attribute |
waitingType | ::mlir::TypeAttr | any type attribute |
barrierStatesType | ::mlir::TypeAttr | any type attribute |
Operands:¶
Operand |
Description |
|---|---|
|
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_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)
Verify that the entry corresponding to the buffer in outstandingCommits tensor is 0.
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_clear_outstanding_commits_transfer_reads (triton::instrument::ExperimentalClearOutstandingCommitsTransferReadsOp)¶
Clear distant outstanding commits and set read visibility for current thread.
Syntax:
operation ::= `tti.experimental_clear_outstanding_commits_transfer_reads` $thread `,` $transferThreadMask `{` $outstandingCommits `(` $outstandingCommitsType `)` `}` `,` $readVisibility `(` $readVisibilityType `)` (`,` $pred^)? attr-dict `:` type($outstandingCommits) `,` type($readVisibility)
For all rows where outstanding commits are more distant than outstandingNum from the current thread,
clear those entries and mark the corresponding buffer as visible in the read visibility table for the threads set in threadMask.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Attributes:¶
| Attribute | MLIR Type | Description |
|---|---|---|
thread | ::mlir::IntegerAttr | 32-bit signless integer attribute |
transferThreadMask | ::mlir::IntegerAttr | 64-bit signless integer attribute |
outstandingNum | ::mlir::IntegerAttr | 32-bit signless integer attribute |
outstandingCommitsType | ::mlir::TypeAttr | any type attribute |
readVisibilityType | ::mlir::TypeAttr | any type attribute |
Operands:¶
Operand |
Description |
|---|---|
|
ptr or ranked tensor of ptr values |
|
ptr or ranked tensor of ptr values |
|
1-bit signless integer |
tti.experimental_clear_outstanding_commits_transfer_writes (triton::instrument::ExperimentalClearOutstandingCommitsTransferWritesOp)¶
Clear distant outstanding commits and set write visibility for current thread.
Syntax:
operation ::= `tti.experimental_clear_outstanding_commits_transfer_writes` $thread `,` $transferThreadMask `{` $outstandingCommits `(` $outstandingCommitsType `)` `}` `,` $writeVisibility `(` $writeVisibilityType `)` (`,` $pred^)? attr-dict `:` type($outstandingCommits) `,` type($writeVisibility)
For all rows where outstanding commits are more distant than outstandingNum from the current thread,
clear those entries and mark the corresponding buffer as visible in the write visibility table for the threads set in threadMask.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Attributes:¶
| Attribute | MLIR Type | Description |
|---|---|---|
thread | ::mlir::IntegerAttr | 32-bit signless integer attribute |
transferThreadMask | ::mlir::IntegerAttr | 64-bit signless integer attribute |
outstandingNum | ::mlir::IntegerAttr | 32-bit signless integer attribute |
outstandingCommitsType | ::mlir::TypeAttr | any type attribute |
writeVisibilityType | ::mlir::TypeAttr | any type attribute |
Operands:¶
Operand |
Description |
|---|---|
|
ptr or ranked tensor of ptr values |
|
ptr or ranked tensor of ptr values |
|
1-bit signless integer |
tti.experimental_clear_read_tracking (triton::instrument::ExperimentalClearReadTrackingOp)¶
Clear the read tracking for a buffer
Syntax:
operation ::= `tti.experimental_clear_read_tracking` $buf `{` $buffers `,` $readTracking `(` $readTrackingType `)` `}` (`,` $pred^)? attr-dict `:` type($buf) `,` type($buffers) `,` type($readTracking)
Clear the read tracking for a buffer.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Attributes:¶
| Attribute | MLIR Type | Description |
|---|---|---|
readTrackingType | ::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_read_visibility (triton::instrument::ExperimentalClearReadVisibilityOp)¶
Clear the read visibility for a buffer
Syntax:
operation ::= `tti.experimental_clear_read_visibility` $buf `{` $buffers `,` $readVisibility `(` $readVisibilityType `)` `}` (`,` $pred^)? attr-dict `:` type($buf) `,` type($buffers) `,` type($readVisibility)
Clear the read visibility for a buffer.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Attributes:¶
| Attribute | MLIR Type | Description |
|---|---|---|
readVisibilityType | ::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_waiting (triton::instrument::ExperimentalClearWaitingOp)¶
Clear the waiting state for the given base thread
Syntax:
operation ::= `tti.experimental_clear_waiting` $mbar `,` $baseThread `{` $barriers `,` $waiting `(` $waitingType `)` `}` (`,` $pred^)? attr-dict `:` type($mbar) `,` type($barriers) `,` type($waiting)
For the barrier row matching mbar, clear both the waiting flag and stored phase for baseThread.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Attributes:¶
| Attribute | MLIR Type | Description |
|---|---|---|
baseThread | ::mlir::IntegerAttr | 32-bit signless integer attribute |
waitingType | ::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_tracking (triton::instrument::ExperimentalClearWriteTrackingOp)¶
Clear the write tracking for a buffer
Syntax:
operation ::= `tti.experimental_clear_write_tracking` $buf `{` $buffers `,` $writeTracking `(` $writeTrackingType `)` `}` (`,` $pred^)? attr-dict `:` type($buf) `,` type($buffers) `,` type($writeTracking)
Clear all the information about threads writing to a buffer.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Attributes:¶
| Attribute | MLIR Type | Description |
|---|---|---|
writeTrackingType | ::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_commit_accesses (triton::instrument::ExperimentalCommitAccessesOp)¶
Commit all the staged accesses for all the buffers.
Syntax:
operation ::= `tti.experimental_commit_accesses` $thread `{` $outstandingCommits `(` $outstandingCommitsType `)` `}` (`,` $pred^)? attr-dict `:` type($outstandingCommits)
Increment the value in outstandingCommits tensor for each entry greater than 0.
Change all the -1 entries in outstandingCommits tensor to 1, signifying 1 outstanding commit.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Attributes:¶
| Attribute | MLIR Type | Description |
|---|---|---|
thread | ::mlir::IntegerAttr | 32-bit signless integer attribute |
outstandingCommitsType | ::mlir::TypeAttr | any type attribute |
Operands:¶
Operand |
Description |
|---|---|
|
ptr or ranked tensor of ptr values |
|
1-bit signless integer |
tti.experimental_copy_read_visibility (triton::instrument::ExperimentalCopyReadVisibilityOp)¶
Replicate read visibility rows from one thread to others
Syntax:
operation ::= `tti.experimental_copy_read_visibility` $sourceThread `,` $destMask `{` $readVisibility `(` $readVisibilityType `)` `}` (`,` $pred^)? attr-dict `:` type($readVisibility)
Copy the read-visibility row of sourceThread to every thread listed in destMask for all buffers. Destination rows are overwritten.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Attributes:¶
| Attribute | MLIR Type | Description |
|---|---|---|
sourceThread | ::mlir::IntegerAttr | 32-bit signless integer attribute |
destMask | ::mlir::IntegerAttr | 64-bit signless integer attribute |
readVisibilityType | ::mlir::TypeAttr | any type attribute |
Operands:¶
Operand |
Description |
|---|---|
|
ptr or ranked tensor of ptr values |
|
1-bit signless integer |
tti.experimental_copy_write_visibility (triton::instrument::ExperimentalCopyWriteVisibilityOp)¶
Replicate write visibility from one thread to others
Syntax:
operation ::= `tti.experimental_copy_write_visibility` $sourceThread `,` $destMask `{` $writeVisibility `(` $writeVisibilityType `)` `}` (`,` $pred^)? attr-dict `:` type($writeVisibility)
Copy the write-visibility bit of sourceThread to every thread listed in destMask for all buffers. Destination bits are overwritten.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Attributes:¶
| Attribute | MLIR Type | Description |
|---|---|---|
sourceThread | ::mlir::IntegerAttr | 32-bit signless integer attribute |
destMask | ::mlir::IntegerAttr | 64-bit signless integer attribute |
writeVisibilityType | ::mlir::TypeAttr | any type attribute |
Operands:¶
Operand |
Description |
|---|---|
|
ptr or ranked tensor of ptr values |
|
1-bit signless integer |
tti.experimental_init_barrier_state (triton::instrument::ExperimentalInitBarrierStateOp)¶
Initialize the auxiliary barrier state
Syntax:
operation ::= `tti.experimental_init_barrier_state` $mbar `,` $count `{` $barriers `,` $states `(` $statesType `)` `}` attr-dict `:` type($mbar) `,` type($barriers) `,` type($states)
Initialize the tracked barrier state to phase 0 and set both the initial and current arrival counts.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Attributes:¶
| Attribute | MLIR Type | Description |
|---|---|---|
count | ::mlir::IntegerAttr | 32-bit signless integer attribute |
statesType | ::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 |
tti.experimental_lock_acquire (triton::instrument::ExperimentalLockAcquireOp)¶
Acquire a lock.
Syntax:
operation ::= `tti.experimental_lock_acquire` $lock (`,` $pred^)? attr-dict `:` type($lock)
Enter a critical section by acquiring a lock with single thread.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Operands:¶
Operand |
Description |
|---|---|
|
ptr or ranked tensor of ptr values |
|
1-bit signless integer |
tti.experimental_lock_release (triton::instrument::ExperimentalLockReleaseOp)¶
Release a lock.
Syntax:
operation ::= `tti.experimental_lock_release` $lock (`,` $pred^)? attr-dict `:` type($lock)
Leave a critical section by releasing a lock with single thread.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Operands:¶
Operand |
Description |
|---|---|
|
ptr or ranked tensor of ptr values |
|
1-bit signless integer |
tti.experimental_set_read_visibility (triton::instrument::ExperimentalSetReadVisibilityOp)¶
Set the read visibility for a buffer
Syntax:
operation ::= `tti.experimental_set_read_visibility` $buf `,` $threadMask `{` $buffers `,` $readVisibility `(` $readVisibilityType `)` `}` (`,` $pred^)? attr-dict `:` type($buf) `,` type($buffers) `,` type($readVisibility)
Add the threads set in threadMask to the read visibility bitmask for the buffer.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Attributes:¶
| Attribute | MLIR Type | Description |
|---|---|---|
threadMask | ::mlir::IntegerAttr | 64-bit signless integer attribute |
readVisibilityType | ::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_set_waiting (triton::instrument::ExperimentalSetWaitingOp)¶
Mark the base thread as waiting on the given barrier phase
Syntax:
operation ::= `tti.experimental_set_waiting` $mbar `,` $baseThread `,` $phase `{` $barriers `,` $waiting `(` $waitingType `)` `}` (`,` $pred^)? attr-dict `:` type($mbar) `,` type($barriers) `,` type($waiting)
For the barrier row matching mbar, set the waiting flag for baseThread and record the barrier phase being waited on.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Attributes:¶
| Attribute | MLIR Type | Description |
|---|---|---|
baseThread | ::mlir::IntegerAttr | 32-bit signless integer attribute |
waitingType | ::mlir::TypeAttr | any type attribute |
Operands:¶
Operand |
Description |
|---|---|
|
memory descriptor type ( |
|
32-bit signless integer |
|
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_visibility (triton::instrument::ExperimentalSetWriteVisibilityOp)¶
Set the write visibility for a buffer
Syntax:
operation ::= `tti.experimental_set_write_visibility` $buf `,` $threadMask `{` $buffers `,` $writeVisibility `(` $writeVisibilityType `)` `}` (`,` $pred^)? attr-dict `:` type($buf) `,` type($buffers) `,` type($writeVisibility)
Set the write visibility for a buffer. Marks the buffer as visible to the threads set in threadMask. Clears out any other threads from the visibility bitmask. We know this is safe because there cannot be outstanding writes to this buffer at this point.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Attributes:¶
| Attribute | MLIR Type | Description |
|---|---|---|
threadMask | ::mlir::IntegerAttr | 64-bit signless integer attribute |
writeVisibilityType | ::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_stage_access_for_commit (triton::instrument::ExperimentalStageAccessForCommitOp)¶
Syntax:
operation ::= `tti.experimental_stage_access_for_commit` $buf `,` $thread `{` $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 |
|---|---|---|
thread | ::mlir::IntegerAttr | 32-bit signless integer attribute |
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 |
tti.experimental_track_visible_reads (triton::instrument::ExperimentalTrackVisibleReadsOp)¶
Start tracking visible reads from a buffer
Syntax:
operation ::= `tti.experimental_track_visible_reads` $mbar `,` $thread `{` $barriers `,` $readVisibility `(` $readVisibilityType `)` `,` $readTracking `(` $readTrackingType `)` `}` (`,` $pred^)? attr-dict `:` type($mbar) `,` type($barriers) `,` type($readVisibility) `,` type($readTracking)
For all buffers that are marked as visible to the thread, start tracking them with a barrier.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Attributes:¶
| Attribute | MLIR Type | Description |
|---|---|---|
thread | ::mlir::IntegerAttr | 32-bit signless integer attribute |
readVisibilityType | ::mlir::TypeAttr | any type attribute |
readTrackingType | ::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_track_visible_writes (triton::instrument::ExperimentalTrackVisibleWritesOp)¶
Start tracking visible writes to a buffer
Syntax:
operation ::= `tti.experimental_track_visible_writes` $mbar `,` $thread `{` $barriers `,` $writeVisibility `(` $writeVisibilityType `)` `,` $writeTracking `(` $writeTrackingType `)` `}` (`,` $pred^)? attr-dict `:` type($mbar) `,` type($barriers) `,` type($writeVisibility) `,` type($writeTracking)
For all buffers that are marked as visible to the thread, start tracking them with a barrier.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Attributes:¶
| Attribute | MLIR Type | Description |
|---|---|---|
thread | ::mlir::IntegerAttr | 32-bit signless integer attribute |
writeVisibilityType | ::mlir::TypeAttr | any type attribute |
writeTrackingType | ::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_transfer_visible_reads (triton::instrument::ExperimentalTransferVisibleReadsOp)¶
Transfer all the reads tracked by a barrier to visible by thread
Syntax:
operation ::= `tti.experimental_transfer_visible_reads` $mbar `,` $threadMask `{` $barriers `,` $readVisibility `(` $readVisibilityType `)` `,` $readTracking `(` $readTrackingType `)` `}` (`,` $pred^)? attr-dict `:` type($mbar) `,` type($barriers) `,` type($readVisibility) `,` type($readTracking)
Transfer all the reads tracked by a barrier to visible by threads set in threadMask.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Attributes:¶
| Attribute | MLIR Type | Description |
|---|---|---|
threadMask | ::mlir::IntegerAttr | 64-bit signless integer attribute |
readVisibilityType | ::mlir::TypeAttr | any type attribute |
readTrackingType | ::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_transfer_visible_writes (triton::instrument::ExperimentalTransferVisibleWritesOp)¶
Transfer all the writes tracked by a barrier to visible by thread
Syntax:
operation ::= `tti.experimental_transfer_visible_writes` $mbar `,` $threadMask `{` $barriers `,` $writeVisibility `(` $writeVisibilityType `)` `,` $writeTracking `(` $writeTrackingType `)` `}` (`,` $pred^)? attr-dict `:` type($mbar) `,` type($barriers) `,` type($writeVisibility) `,` type($writeTracking)
Transfer all the writes tracked by a barrier to visible by threads set in threadMask.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Attributes:¶
| Attribute | MLIR Type | Description |
|---|---|---|
threadMask | ::mlir::IntegerAttr | 64-bit signless integer attribute |
writeVisibilityType | ::mlir::TypeAttr | any type attribute |
writeTrackingType | ::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_update_barrier_state (triton::instrument::ExperimentalUpdateBarrierStateOp)¶
Update the auxiliary barrier state after a verified arrive
Syntax:
operation ::= `tti.experimental_update_barrier_state` $mbar `,` $count `{` $barriers `,` $states `(` $statesType `)` `}` (`,` $pred^)? attr-dict `:` type($mbar) `,` type($barriers) `,` type($states)
Apply an arrive count to the tracked barrier state, toggling the phase when the count reaches zero and reloading the current count from the initial count.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Attributes:¶
| Attribute | MLIR Type | Description |
|---|---|---|
count | ::mlir::IntegerAttr | 32-bit signless integer attribute |
statesType | ::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_verify_barrier_arrive (triton::instrument::ExperimentalVerifyBarrierArriveOp)¶
Verify an arrive count against the tracked barrier state
Syntax:
operation ::= `tti.experimental_verify_barrier_arrive` $mbar `,` $count `{` $barriers `,` $states `(` $statesType `)` `}` (`,` $pred^)? attr-dict `:` type($mbar) `,` type($barriers) `,` type($states)
Check that applying the arrive count would not drive the tracked current count negative. Triggers an assertion on failure.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Attributes:¶
| Attribute | MLIR Type | Description |
|---|---|---|
count | ::mlir::IntegerAttr | 32-bit signless integer attribute |
statesType | ::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_verify_read_visibility (triton::instrument::ExperimentalVerifyReadVisibilityOp)¶
Verify that all the reads from a buffer are visible to the thread
Syntax:
operation ::= `tti.experimental_verify_read_visibility` $buf `,` $thread `{` $buffers `,` $readVisibility `(` $readVisibilityType `)` `}` `,` $operandName (`,` $pred^)? attr-dict `:` type($buf) `,` type($buffers) `,` type($readVisibility)
Verify that all the reads from a buffer are visible to the thread.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Attributes:¶
| Attribute | MLIR Type | Description |
|---|---|---|
thread | ::mlir::IntegerAttr | 32-bit signless integer attribute |
readVisibilityType | ::mlir::TypeAttr | any type attribute |
operandName | ::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_verify_write_visibility (triton::instrument::ExperimentalVerifyWriteVisibilityOp)¶
Verify that the buffer write is visible to the thread
Syntax:
operation ::= `tti.experimental_verify_write_visibility` $buf `,` $thread `{` $buffers `,` $writeVisibility `(` $writeVisibilityType `)` `}` `,` $operandName (`,` $pred^)? attr-dict `:` type($buf) `,` type($buffers) `,` type($writeVisibility)
Verify that the buffer write is visible to the thread, or no threads are writing to the buffer.
Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
Attributes:¶
| Attribute | MLIR Type | Description |
|---|---|---|
thread | ::mlir::IntegerAttr | 32-bit signless integer attribute |
writeVisibilityType | ::mlir::TypeAttr | any type attribute |
operandName | ::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 |