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布尔属性

操作数:

操作数

描述

condition

1位无符号整数值的张量

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

检查由 mbar 保护的缓冲区是否有未完成的读取操作

语法

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

检查由 mbar 保护的缓冲区是否有未完成的读取操作。

接口：MemoryEffectOpInterface (MemoryEffectOpInterface)

效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}

属性：

属性MLIR 类型描述
readBarsType::mlir::TypeAttr任何类型属性

操作数:

操作数

描述

buf

内存描述符类型 (::mlir::triton::gpu::MemDescType) 在 Triton IR 类型系统中

buffers

浮点数、整数或指针值的有秩张量

readBars

指针或指针值的有秩张量

pred

1位无符号整数

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

检查由 mbar 保护的缓冲区是否有未完成的写入操作

语法

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

检查由 mbar 保护的缓冲区是否有未完成的写入操作。

接口：MemoryEffectOpInterface (MemoryEffectOpInterface)

效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}

属性：

属性MLIR 类型描述
writeBarsType::mlir::TypeAttr任何类型属性

操作数:

操作数

描述

buf

内存描述符类型 (::mlir::triton::gpu::MemDescType) 在 Triton IR 类型系统中

buffers

浮点数、整数或指针值的有秩张量

writeBars

指针或指针值的有秩张量

pred

1位无符号整数

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

检查缓冲区是否有未完成的写入提交。

语法

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

检查缓冲区是否有未完成的写入提交。

接口：MemoryEffectOpInterface (MemoryEffectOpInterface)

效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}

属性：

属性MLIR 类型描述
writeCommitsType::mlir::TypeAttr任何类型属性

操作数：

操作数

描述

buf

内存描述符类型 (::mlir::triton::gpu::MemDescType) 在 Triton IR 类型系统中

buffers

浮点数、整数或指针值的有秩张量

writeCommits

指针或指针值的有秩张量

pred

1位无符号整数

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

清除由 mbar 保护的缓冲区的读取状态

语法

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

清除由 mbar 保护的缓冲区的读取状态。

接口：MemoryEffectOpInterface (MemoryEffectOpInterface)

效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}

属性:

属性MLIR 类型描述
readBarsType::mlir::TypeAttr任何类型属性

操作数:

操作数

描述

mbar

内存描述符类型 (::mlir::triton::gpu::MemDescType) 在 Triton IR 类型系统中

barriers

浮点数、整数或指针值的有秩张量

readBars

指针或指针值的有秩张量

pred

1位无符号整数

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

清除由 mbar 保护的缓冲区的写入状态

语法

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

清除由 mbar 保护的缓冲区的写入状态。

接口：MemoryEffectOpInterface (MemoryEffectOpInterface)

效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}

属性：

属性MLIR 类型描述
writeBarsType::mlir::TypeAttr任何类型属性

操作数:

操作数

描述

mbar

内存描述符类型 (::mlir::triton::gpu::MemDescType) 在 Triton IR 类型系统中

writeBars

指针或指针值的有秩张量

pred

1位无符号整数

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

清除所有距离超过 outstandingNum 的写入提交。

语法

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

清除所有距离当前线程超过 outstandingNum 的写入提交。

接口：MemoryEffectOpInterface (MemoryEffectOpInterface)

效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}

属性:

属性MLIR 类型描述
writeCommitsType::mlir::TypeAttr任何类型属性
outstandingNum::mlir::IntegerAttr32位无符号整数属性

操作数：

操作数

描述

writeCommits

指针或指针值的有秩张量

pred

1位无符号整数

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

提交所有缓冲区的暂存写入。

语法

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

提交所有缓冲区的暂存写入。

接口：MemoryEffectOpInterface (MemoryEffectOpInterface)

效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}

属性：

属性MLIR 类型描述
writeCommitsType::mlir::TypeAttr任何类型属性

操作数:

操作数

描述

writeCommits

指针或指针值的有秩张量

pred

1位无符号整数

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

将缓冲区标记为通过 mbar 作为守卫进行读取

语法

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)

将缓冲区标记为通过 mbar 作为守卫进行读取。

接口：MemoryEffectOpInterface (MemoryEffectOpInterface)

效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}

属性:

属性MLIR 类型描述
readBarsType::mlir::TypeAttr任何类型属性

操作数：

操作数

描述

buf

内存描述符类型 (::mlir::triton::gpu::MemDescType) 在 Triton IR 类型系统中

mbar

内存描述符类型 (::mlir::triton::gpu::MemDescType) 在 Triton IR 类型系统中

buffers

浮点数、整数或指针值的有秩张量

barriers

浮点数、整数或指针值的有秩张量

readBars

指针或指针值的有秩张量

pred

1位无符号整数

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

将缓冲区标记为通过 mbar 作为守卫进行写入

语法

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

将缓冲区标记为通过 mbar 作为守卫进行写入。

接口：MemoryEffectOpInterface (MemoryEffectOpInterface)

效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}

属性：

属性MLIR 类型描述
writeBarsType::mlir::TypeAttr任何类型属性

操作数：

操作数

描述

buf

内存描述符类型 (::mlir::triton::gpu::MemDescType) 在 Triton IR 类型系统中

mbar

内存描述符类型 (::mlir::triton::gpu::MemDescType) 在 Triton IR 类型系统中

buffers

浮点数、整数或指针值的有秩张量

writeBars

指针或指针值的有秩张量

pred

1位无符号整数

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

定义指向共享内存缓冲区的指针数组

语法

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

创建一个指向共享内存缓冲区的指针张量。

特性: AlwaysSpeculatableImplTrait

接口: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

效果: MemoryEffects::Effect{}

属性：

属性MLIR 类型描述
offsets::mlir::DenseI32ArrayAttri32 密集数组属性

结果:

结果

描述

result

浮点数、整数或指针值的有秩张量

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

准备进行缓冲区的异步复制。暂存直至调用 commit_group。

语法

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

准备进行缓冲区的异步复制。暂存直至调用 commit_group。实现会将 -1 写入到 write_commits 张量中与缓冲区对应的索引下。

接口：MemoryEffectOpInterface (MemoryEffectOpInterface)

效果：MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}

属性：

属性MLIR 类型描述
writeCommitsType::mlir::TypeAttr任何类型属性

操作数：

操作数

描述

buf

内存描述符类型 (::mlir::triton::gpu::MemDescType) 在 Triton IR 类型系统中

buffers

浮点数、整数或指针值的有秩张量

writeCommits

指针或指针值的有秩张量

pred

1位无符号整数