TritonInstrumentOps¶
tti.experimental_assert_in_thread (triton::instrument::ExperimentalAssertInThreadOp)¶
在当前线程内断言条件
语法
operation ::= `tti.experimental_assert_in_thread` $condition `,` $message attr-dict `:` type($condition)
断言在当前线程中所有值都可用的情况下，该条件为真。如果条件为假，则打印消息并中止程序。如果 check_any 为真，则条件中的任何一个值必须为真。否则，条件中的所有值都必须为真。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
message | ::mlir::StringAttr | 字符串属性 |
check_any | ::mlir::BoolAttr | 布尔属性 |
操作数:¶
操作数 |
描述 |
|---|---|
|
1 位无符号整数值的张量 |
tti.experimental_buffer_pointers (triton::instrument::ExperimentalBufferPointersOp)¶
定义一个指向共享内存缓冲区的指针数组
语法
operation ::= `tti.experimental_buffer_pointers` $offsets `,` $memType attr-dict `:` type($result)
创建一个指向共享内存缓冲区的指针张量。
特性: AlwaysSpeculatableImplTrait
接口: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)
效果: MemoryEffects::Effect{}
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
offsets | ::mlir::DenseI32ArrayAttr | i32 密集数组属性 |
memType | ::mlir::triton::instrument::MemTypeAttr | 允许的32位无符号整数情况：0, 1 |
结果:¶
结果 |
描述 |
|---|---|
|
浮点数、整数或 ptr 值的秩张量 |
tti.experimental_check_all_active_waiting (triton::instrument::ExperimentalCheckAllActiveWaitingOp)¶
断言并非所有活动线程都在等待匹配的阶段
语法
operation ::= `tti.experimental_check_all_active_waiting` $activeMask `,` $barriers `,` $waiting `(` $waitingType `)` `,` $barrierStates `(` $barrierStatesType `)` (`,` $pred^)? attr-dict `:` type($barriers) `,` type($waiting) `,` type($barrierStates)
将等待线程筛选为记录的阶段与当前屏障阶段匹配的线程，跨屏障进行“或”归约，并断言 (waitingMask & activeMask) != activeMask。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
activeMask | ::mlir::IntegerAttr | 32位无符号整数属性 |
waitingType | ::mlir::TypeAttr | 任意类型属性 |
barrierStatesType | ::mlir::TypeAttr | 任意类型属性 |
操作数：¶
操作数 |
描述 |
|---|---|
|
浮点数、整数或 ptr 值的秩张量 |
|
ptr 或 ptr 值的秩张量 |
|
ptr 或 ptr 值的秩张量 |
|
1位无符号整数 |
tti.experimental_check_outstanding_commits (triton::instrument::ExperimentalCheckOutstandingCommitsOp)¶
检查缓冲区是否有未完成的提交。
语法
operation ::= `tti.experimental_check_outstanding_commits` $buf `{` $buffers `,` $outstandingCommits `(` $outstandingCommitsType `)` `}` `,` $pendingAccessType (`,` $pred^)? attr-dict `:` type($buf) `,` type($buffers) `,` type($outstandingCommits)
验证 outstandingCommits 张量中与缓冲区对应的条目是否为 0。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
outstandingCommitsType | ::mlir::TypeAttr | 任意类型属性 |
pendingAccessType | ::mlir::StringAttr | 字符串属性 |
操作数:¶
操作数 |
描述 |
|---|---|
|
Triton IR 类型系统中的内存描述符类型 ( |
|
浮点数、整数或 ptr 值的秩张量 |
|
ptr 或 ptr 值的秩张量 |
|
1位无符号整数 |
tti.experimental_clear_outstanding_commits_transfer_reads (triton::instrument::ExperimentalClearOutstandingCommitsTransferReadsOp)¶
清除远端的未完成提交，并为当前线程设置读取可见性。
语法
operation ::= `tti.experimental_clear_outstanding_commits_transfer_reads` $thread `,` $transferThreadMask `{` $outstandingCommits `(` $outstandingCommitsType `)` `}` `,` $readVisibility `(` $readVisibilityType `)` (`,` $pred^)? attr-dict `:` type($outstandingCommits) `,` type($readVisibility)
对于所有未完成提交距离当前线程比 outstandingNum 更远的行，清除这些条目，并在读取可见性表中将相应的缓冲区标记为对 threadMask 中设置的线程可见。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
thread | ::mlir::IntegerAttr | 32位无符号整数属性 |
transferThreadMask | ::mlir::IntegerAttr | 64 位无符号整数属性 |
outstandingNum | ::mlir::IntegerAttr | 32位无符号整数属性 |
outstandingCommitsType | ::mlir::TypeAttr | 任意类型属性 |
readVisibilityType | ::mlir::TypeAttr | 任意类型属性 |
操作数：¶
操作数 |
描述 |
|---|---|
|
ptr 或 ptr 值的秩张量 |
|
ptr 或 ptr 值的秩张量 |
|
1位无符号整数 |
tti.experimental_clear_outstanding_commits_transfer_writes (triton::instrument::ExperimentalClearOutstandingCommitsTransferWritesOp)¶
清除远端的未完成提交，并为当前线程设置写入可见性。
语法
operation ::= `tti.experimental_clear_outstanding_commits_transfer_writes` $thread `,` $transferThreadMask `{` $outstandingCommits `(` $outstandingCommitsType `)` `}` `,` $writeVisibility `(` $writeVisibilityType `)` (`,` $pred^)? attr-dict `:` type($outstandingCommits) `,` type($writeVisibility)
对于所有未完成提交距离当前线程比 outstandingNum 更远的行，清除这些条目，并在写入可见性表中将相应的缓冲区标记为对 threadMask 中设置的线程可见。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
thread | ::mlir::IntegerAttr | 32位无符号整数属性 |
transferThreadMask | ::mlir::IntegerAttr | 64 位无符号整数属性 |
outstandingNum | ::mlir::IntegerAttr | 32位无符号整数属性 |
outstandingCommitsType | ::mlir::TypeAttr | 任意类型属性 |
writeVisibilityType | ::mlir::TypeAttr | 任意类型属性 |
操作数:¶
操作数 |
描述 |
|---|---|
|
ptr 或 ptr 值的秩张量 |
|
ptr 或 ptr 值的秩张量 |
|
1位无符号整数 |
tti.experimental_clear_read_tracking (triton::instrument::ExperimentalClearReadTrackingOp)¶
清除缓冲区的读取跟踪
语法
operation ::= `tti.experimental_clear_read_tracking` $buf `{` $buffers `,` $readTracking `(` $readTrackingType `)` `}` (`,` $pred^)? attr-dict `:` type($buf) `,` type($buffers) `,` type($readTracking)
清除缓冲区的读取跟踪。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
readTrackingType | ::mlir::TypeAttr | 任意类型属性 |
操作数：¶
操作数 |
描述 |
|---|---|
|
Triton IR 类型系统中的内存描述符类型 ( |
|
浮点数、整数或 ptr 值的秩张量 |
|
ptr 或 ptr 值的秩张量 |
|
1位无符号整数 |
tti.experimental_clear_read_visibility (triton::instrument::ExperimentalClearReadVisibilityOp)¶
清除缓冲区的读取可见性
语法
operation ::= `tti.experimental_clear_read_visibility` $buf `{` $buffers `,` $readVisibility `(` $readVisibilityType `)` `}` (`,` $pred^)? attr-dict `:` type($buf) `,` type($buffers) `,` type($readVisibility)
清除缓冲区的读取可见性。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
readVisibilityType | ::mlir::TypeAttr | 任意类型属性 |
操作数：¶
操作数 |
描述 |
|---|---|
|
Triton IR 类型系统中的内存描述符类型 ( |
|
浮点数、整数或 ptr 值的秩张量 |
|
ptr 或 ptr 值的秩张量 |
|
1位无符号整数 |
tti.experimental_clear_waiting (triton::instrument::ExperimentalClearWaitingOp)¶
清除给定基础线程的等待状态
语法
operation ::= `tti.experimental_clear_waiting` $mbar `,` $baseThread `{` $barriers `,` $waiting `(` $waitingType `)` `}` (`,` $pred^)? attr-dict `:` type($mbar) `,` type($barriers) `,` type($waiting)
对于与 mbar 匹配的屏障行，清除 baseThread 的等待标志和存储的阶段。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
baseThread | ::mlir::IntegerAttr | 32位无符号整数属性 |
waitingType | ::mlir::TypeAttr | 任意类型属性 |
操作数：¶
操作数 |
描述 |
|---|---|
|
Triton IR 类型系统中的内存描述符类型 ( |
|
浮点数、整数或 ptr 值的秩张量 |
|
ptr 或 ptr 值的秩张量 |
|
1位无符号整数 |
tti.experimental_clear_write_tracking (triton::instrument::ExperimentalClearWriteTrackingOp)¶
清除缓冲区的写入跟踪
语法
operation ::= `tti.experimental_clear_write_tracking` $buf `{` $buffers `,` $writeTracking `(` $writeTrackingType `)` `}` (`,` $pred^)? attr-dict `:` type($buf) `,` type($buffers) `,` type($writeTracking)
清除有关写入缓冲区的所有线程的信息。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
writeTrackingType | ::mlir::TypeAttr | 任意类型属性 |
操作数:¶
操作数 |
描述 |
|---|---|
|
Triton IR 类型系统中的内存描述符类型 ( |
|
浮点数、整数或 ptr 值的秩张量 |
|
ptr 或 ptr 值的秩张量 |
|
1位无符号整数 |
tti.experimental_commit_accesses (triton::instrument::ExperimentalCommitAccessesOp)¶
提交所有缓冲区的全部暂存访问。
语法
operation ::= `tti.experimental_commit_accesses` $thread `{` $outstandingCommits `(` $outstandingCommitsType `)` `}` (`,` $pred^)? attr-dict `:` type($outstandingCommits)
对于每个大于 0 的条目，将 outstandingCommits 张量中的值递增。将 outstandingCommits 张量中所有 -1 的条目更改为 1，表示 1 个未完成的提交。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
thread | ::mlir::IntegerAttr | 32位无符号整数属性 |
outstandingCommitsType | ::mlir::TypeAttr | 任意类型属性 |
操作数 (Operands):¶
操作数 |
描述 |
|---|---|
|
ptr 或 ptr 值的秩张量 |
|
1位无符号整数 |
tti.experimental_copy_read_visibility (triton::instrument::ExperimentalCopyReadVisibilityOp)¶
将一个线程的读取可见性行复制到其他线程
语法
operation ::= `tti.experimental_copy_read_visibility` $sourceThread `,` $destMask `{` $readVisibility `(` $readVisibilityType `)` `}` (`,` $pred^)? attr-dict `:` type($readVisibility)
将 sourceThread 的读取可见性行复制到 destMask 中列出的所有缓冲区的每个线程。目标行将被覆盖。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
sourceThread | ::mlir::IntegerAttr | 32位无符号整数属性 |
destMask | ::mlir::IntegerAttr | 64 位无符号整数属性 |
readVisibilityType | ::mlir::TypeAttr | 任意类型属性 |
操作数：¶
操作数 |
描述 |
|---|---|
|
ptr 或 ptr 值的秩张量 |
|
1位无符号整数 |
tti.experimental_copy_write_visibility (triton::instrument::ExperimentalCopyWriteVisibilityOp)¶
将写入可见性从一个线程复制到其他线程
语法
operation ::= `tti.experimental_copy_write_visibility` $sourceThread `,` $destMask `{` $writeVisibility `(` $writeVisibilityType `)` `}` (`,` $pred^)? attr-dict `:` type($writeVisibility)
将 sourceThread 的写入可见性位复制到 destMask 中列出的所有缓冲区的每个线程。目标位将被覆盖。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
sourceThread | ::mlir::IntegerAttr | 32位无符号整数属性 |
destMask | ::mlir::IntegerAttr | 64 位无符号整数属性 |
writeVisibilityType | ::mlir::TypeAttr | 任意类型属性 |
操作数：¶
操作数 |
描述 |
|---|---|
|
ptr 或 ptr 值的秩张量 |
|
1位无符号整数 |
tti.experimental_init_barrier_state (triton::instrument::ExperimentalInitBarrierStateOp)¶
初始化辅助屏障状态
语法
operation ::= `tti.experimental_init_barrier_state` $mbar `,` $count `{` $barriers `,` $states `(` $statesType `)` `}` attr-dict `:` type($mbar) `,` type($barriers) `,` type($states)
将跟踪的屏障状态初始化为阶段 0，并设置初始和当前的到达计数。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
count | ::mlir::IntegerAttr | 32位无符号整数属性 |
statesType | ::mlir::TypeAttr | 任意类型属性 |
操作数：¶
操作数 |
描述 |
|---|---|
|
Triton IR 类型系统中的内存描述符类型 ( |
|
浮点数、整数或 ptr 值的秩张量 |
|
ptr 或 ptr 值的秩张量 |
tti.experimental_lock_acquire (triton::instrument::ExperimentalLockAcquireOp)¶
获取锁。
语法
operation ::= `tti.experimental_lock_acquire` $lock (`,` $pred^)? attr-dict `:` type($lock)
通过单个线程获取锁来进入临界区。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
操作数：¶
操作数 |
描述 |
|---|---|
|
ptr 或 ptr 值的秩张量 |
|
1位无符号整数 |
tti.experimental_lock_release (triton::instrument::ExperimentalLockReleaseOp)¶
释放锁。
语法
operation ::= `tti.experimental_lock_release` $lock (`,` $pred^)? attr-dict `:` type($lock)
通过单个线程释放锁来离开临界区。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
操作数 (Operands):¶
操作数 |
描述 |
|---|---|
|
ptr 或 ptr 值的秩张量 |
|
1位无符号整数 |
tti.experimental_set_read_visibility (triton::instrument::ExperimentalSetReadVisibilityOp)¶
设置缓冲区的读取可见性
语法
operation ::= `tti.experimental_set_read_visibility` $buf `,` $threadMask `{` $buffers `,` $readVisibility `(` $readVisibilityType `)` `}` (`,` $pred^)? attr-dict `:` type($buf) `,` type($buffers) `,` type($readVisibility)
将 threadMask 中设置的线程添加到缓冲区的读取可见性位掩码中。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
threadMask | ::mlir::IntegerAttr | 64 位无符号整数属性 |
readVisibilityType | ::mlir::TypeAttr | 任意类型属性 |
操作数：¶
操作数 |
描述 |
|---|---|
|
Triton IR 类型系统中的内存描述符类型 ( |
|
浮点数、整数或 ptr 值的秩张量 |
|
ptr 或 ptr 值的秩张量 |
|
1位无符号整数 |
tti.experimental_set_waiting (triton::instrument::ExperimentalSetWaitingOp)¶
将基础线程标记为在给定的屏障阶段等待
语法
operation ::= `tti.experimental_set_waiting` $mbar `,` $baseThread `,` $phase `{` $barriers `,` $waiting `(` $waitingType `)` `}` (`,` $pred^)? attr-dict `:` type($mbar) `,` type($barriers) `,` type($waiting)
对于与 mbar 匹配的屏障行，为 baseThread 设置等待标志并记录正在等待的屏障阶段。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
baseThread | ::mlir::IntegerAttr | 32位无符号整数属性 |
waitingType | ::mlir::TypeAttr | 任意类型属性 |
操作数：¶
操作数 |
描述 |
|---|---|
|
Triton IR 类型系统中的内存描述符类型 ( |
|
32位无符号整数 |
|
浮点数、整数或 ptr 值的秩张量 |
|
ptr 或 ptr 值的秩张量 |
|
1位无符号整数 |
tti.experimental_set_write_visibility (triton::instrument::ExperimentalSetWriteVisibilityOp)¶
设置缓冲区的写入可见性
语法
operation ::= `tti.experimental_set_write_visibility` $buf `,` $threadMask `{` $buffers `,` $writeVisibility `(` $writeVisibilityType `)` `}` (`,` $pred^)? attr-dict `:` type($buf) `,` type($buffers) `,` type($writeVisibility)
设置缓冲区的写入可见性。将缓冲区标记为对 threadMask 中设置的线程可见。从可见性位掩码中清除任何其他线程。我们知道这样做是安全的，因为此时不可能有对该缓冲区的未完成写入。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
threadMask | ::mlir::IntegerAttr | 64 位无符号整数属性 |
writeVisibilityType | ::mlir::TypeAttr | 任意类型属性 |
操作数：¶
操作数 |
描述 |
|---|---|
|
Triton IR 类型系统中的内存描述符类型 ( |
|
浮点数、整数或 ptr 值的秩张量 |
|
ptr 或 ptr 值的秩张量 |
|
1位无符号整数 |
tti.experimental_stage_access_for_commit (triton::instrument::ExperimentalStageAccessForCommitOp)¶
语法
operation ::= `tti.experimental_stage_access_for_commit` $buf `,` $thread `{` $buffers `,` $outstandingCommits `(` $outstandingCommitsType `)` `}` (`,` $pred^)? attr-dict `:` type($buf) `,` type($buffers) `,` type($outstandingCommits)
对于使用 outstanding 来跟踪未完成提交数量（而不是 mbarriers）的操作，通过将其标记为 -1，将缓冲区标记为正在被访问但尚未提交。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
thread | ::mlir::IntegerAttr | 32位无符号整数属性 |
outstandingCommitsType | ::mlir::TypeAttr | 任意类型属性 |
操作数：¶
操作数 |
描述 |
|---|---|
|
Triton IR 类型系统中的内存描述符类型 ( |
|
浮点数、整数或 ptr 值的秩张量 |
|
ptr 或 ptr 值的秩张量 |
|
1位无符号整数 |
tti.experimental_track_visible_reads (triton::instrument::ExperimentalTrackVisibleReadsOp)¶
开始跟踪来自缓冲区的可见读取
语法
operation ::= `tti.experimental_track_visible_reads` $mbar `,` $thread `{` $barriers `,` $readVisibility `(` $readVisibilityType `)` `,` $readTracking `(` $readTrackingType `)` `}` (`,` $pred^)? attr-dict `:` type($mbar) `,` type($barriers) `,` type($readVisibility) `,` type($readTracking)
对于所有标记为对线程可见的缓冲区，开始使用屏障对其进行跟踪。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
thread | ::mlir::IntegerAttr | 32位无符号整数属性 |
readVisibilityType | ::mlir::TypeAttr | 任意类型属性 |
readTrackingType | ::mlir::TypeAttr | 任意类型属性 |
操作数：¶
操作数 |
描述 |
|---|---|
|
Triton IR 类型系统中的内存描述符类型 ( |
|
浮点数、整数或 ptr 值的秩张量 |
|
ptr 或 ptr 值的秩张量 |
|
ptr 或 ptr 值的秩张量 |
|
1位无符号整数 |
tti.experimental_track_visible_writes (triton::instrument::ExperimentalTrackVisibleWritesOp)¶
开始跟踪对缓冲区的可见写入
语法
operation ::= `tti.experimental_track_visible_writes` $mbar `,` $thread `{` $barriers `,` $writeVisibility `(` $writeVisibilityType `)` `,` $writeTracking `(` $writeTrackingType `)` `}` (`,` $pred^)? attr-dict `:` type($mbar) `,` type($barriers) `,` type($writeVisibility) `,` type($writeTracking)
对于所有标记为对线程可见的缓冲区，开始使用屏障对其进行跟踪。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
thread | ::mlir::IntegerAttr | 32位无符号整数属性 |
writeVisibilityType | ::mlir::TypeAttr | 任意类型属性 |
writeTrackingType | ::mlir::TypeAttr | 任意类型属性 |
操作数：¶
操作数 |
描述 |
|---|---|
|
Triton IR 类型系统中的内存描述符类型 ( |
|
浮点数、整数或 ptr 值的秩张量 |
|
ptr 或 ptr 值的秩张量 |
|
ptr 或 ptr 值的秩张量 |
|
1位无符号整数 |
tti.experimental_transfer_visible_reads (triton::instrument::ExperimentalTransferVisibleReadsOp)¶
将屏障跟踪的所有读取转移为线程可见
语法
operation ::= `tti.experimental_transfer_visible_reads` $mbar `,` $threadMask `{` $barriers `,` $readVisibility `(` $readVisibilityType `)` `,` $readTracking `(` $readTrackingType `)` `}` (`,` $pred^)? attr-dict `:` type($mbar) `,` type($barriers) `,` type($readVisibility) `,` type($readTracking)
将屏障跟踪的所有读取转移为 threadMask 中设置的线程可见。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
threadMask | ::mlir::IntegerAttr | 64 位无符号整数属性 |
readVisibilityType | ::mlir::TypeAttr | 任意类型属性 |
readTrackingType | ::mlir::TypeAttr | 任意类型属性 |
操作数：¶
操作数 |
描述 |
|---|---|
|
Triton IR 类型系统中的内存描述符类型 ( |
|
浮点数、整数或 ptr 值的秩张量 |
|
ptr 或 ptr 值的秩张量 |
|
ptr 或 ptr 值的秩张量 |
|
1位无符号整数 |
tti.experimental_transfer_visible_writes (triton::instrument::ExperimentalTransferVisibleWritesOp)¶
将屏障跟踪的所有写入转移为线程可见
语法
operation ::= `tti.experimental_transfer_visible_writes` $mbar `,` $threadMask `{` $barriers `,` $writeVisibility `(` $writeVisibilityType `)` `,` $writeTracking `(` $writeTrackingType `)` `}` (`,` $pred^)? attr-dict `:` type($mbar) `,` type($barriers) `,` type($writeVisibility) `,` type($writeTracking)
将屏障跟踪的所有写入转移为 threadMask 中设置的线程可见。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
threadMask | ::mlir::IntegerAttr | 64 位无符号整数属性 |
writeVisibilityType | ::mlir::TypeAttr | 任意类型属性 |
writeTrackingType | ::mlir::TypeAttr | 任意类型属性 |
操作数：¶
操作数 |
描述 |
|---|---|
|
Triton IR 类型系统中的内存描述符类型 ( |
|
浮点数、整数或 ptr 值的秩张量 |
|
ptr 或 ptr 值的秩张量 |
|
ptr 或 ptr 值的秩张量 |
|
1位无符号整数 |
tti.experimental_update_barrier_state (triton::instrument::ExperimentalUpdateBarrierStateOp)¶
在经过验证的到达后更新辅助屏障状态
语法
operation ::= `tti.experimental_update_barrier_state` $mbar `,` $count `{` $barriers `,` $states `(` $statesType `)` `}` (`,` $pred^)? attr-dict `:` type($mbar) `,` type($barriers) `,` type($states)
将一个到达计数应用于跟踪的屏障状态，当计数达到零时切换阶段，并从初始计数重新加载当前计数。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
count | ::mlir::IntegerAttr | 32位无符号整数属性 |
statesType | ::mlir::TypeAttr | 任意类型属性 |
操作数：¶
操作数 |
描述 |
|---|---|
|
Triton IR 类型系统中的内存描述符类型 ( |
|
浮点数、整数或 ptr 值的秩张量 |
|
ptr 或 ptr 值的秩张量 |
|
1位无符号整数 |
tti.experimental_verify_barrier_arrive (triton::instrument::ExperimentalVerifyBarrierArriveOp)¶
根据跟踪的屏障状态验证到达计数
语法
operation ::= `tti.experimental_verify_barrier_arrive` $mbar `,` $count `{` $barriers `,` $states `(` $statesType `)` `}` (`,` $pred^)? attr-dict `:` type($mbar) `,` type($barriers) `,` type($states)
检查应用到达计数不会使跟踪的当前计数变为负数。失败时触发断言。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
count | ::mlir::IntegerAttr | 32位无符号整数属性 |
statesType | ::mlir::TypeAttr | 任意类型属性 |
操作数：¶
操作数 |
描述 |
|---|---|
|
Triton IR 类型系统中的内存描述符类型 ( |
|
浮点数、整数或 ptr 值的秩张量 |
|
ptr 或 ptr 值的秩张量 |
|
1位无符号整数 |
tti.experimental_verify_read_visibility (triton::instrument::ExperimentalVerifyReadVisibilityOp)¶
验证来自缓冲区的所有读取对线程都是可见的
语法
operation ::= `tti.experimental_verify_read_visibility` $buf `,` $thread `{` $buffers `,` $readVisibility `(` $readVisibilityType `)` `}` `,` $operandName (`,` $pred^)? attr-dict `:` type($buf) `,` type($buffers) `,` type($readVisibility)
验证来自缓冲区的所有读取对线程都是可见的。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
thread | ::mlir::IntegerAttr | 32位无符号整数属性 |
readVisibilityType | ::mlir::TypeAttr | 任意类型属性 |
operandName | ::mlir::StringAttr | 字符串属性 |
操作数：¶
操作数 |
描述 |
|---|---|
|
Triton IR 类型系统中的内存描述符类型 ( |
|
浮点数、整数或 ptr 值的秩张量 |
|
ptr 或 ptr 值的秩张量 |
|
1位无符号整数 |
tti.experimental_verify_write_visibility (triton::instrument::ExperimentalVerifyWriteVisibilityOp)¶
验证缓冲区写入对线程是可见的
语法
operation ::= `tti.experimental_verify_write_visibility` $buf `,` $thread `{` $buffers `,` $writeVisibility `(` $writeVisibilityType `)` `}` `,` $operandName (`,` $pred^)? attr-dict `:` type($buf) `,` type($buffers) `,` type($writeVisibility)
验证缓冲区写入对线程是可见的，或者没有线程正在写入该缓冲区。
接口：MemoryEffectOpInterface (MemoryEffectOpInterface)
效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
thread | ::mlir::IntegerAttr | 32位无符号整数属性 |
writeVisibilityType | ::mlir::TypeAttr | 任意类型属性 |
operandName | ::mlir::StringAttr | 字符串属性 |
操作数：¶
操作数 |
描述 |
|---|---|
|
Triton IR 类型系统中的内存描述符类型 ( |
|
浮点数、整数或 ptr 值的秩张量 |
|
ptr 或 ptr 值的秩张量 |
|
1位无符号整数 |