TritonInstrumentOps

`tti.experimental_assert_uniform` (triton::instrument::ExperimentalAssertUniformOp)

断言一致性条件

语法

operation ::= `tti.experimental_assert_uniform` $condition `,` $message attr-dict-with-keyword

断言该条件为真，前提是线程束组中的所有线程都具有相同的值，因此只需一个线程即可评估断言并打印消息。

接口：MemoryEffectOpInterface (MemoryEffectOpInterface)

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

属性：

属性	MLIR 类型	描述
`消息`	::mlir::StringAttr	字符串属性

操作数：

操作数	描述
`条件`	1位无符号整数

`tti.experimental_buffer_descriptors` (triton::instrument::ExperimentalBufferDescriptorsOp)

定义缓冲区描述符数组

语法

operation ::= `tti.experimental_buffer_descriptors` $offsets `,` $lengths `,` $memType attr-dict `:` type($result)

创建一个缓冲区描述符张量，将 32 位指针偏移量和 32 位长度打包到 64 位元素中。

特性: AlwaysSpeculatableImplTrait

接口: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

效果: MemoryEffects::Effect{}

属性：

属性	MLIR 类型	描述
`offsets`	::mlir::DenseI32ArrayAttr	i32 密集数组属性
`长度`	::mlir::DenseI32ArrayAttr	i32 密集数组属性
`内存类型`	::mlir::triton::instrument::MemTypeAttr	允许的32位无符号整数情况：0, 1

结果：

结果	描述
`result`	浮点数、整数或 ptr 值的秩张量

`tti.experimental_cluster_cta_id` (triton::instrument::ExperimentalClusterCTAIdOp)

获取当前集群内的 CTA ID

语法

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

返回用于索引 ConSan 多 CTA 暂存块（scratch slab）的集群局部 CTA ID。对于单 CTA 内核，此值始终为零。

特性: AlwaysSpeculatableImplTrait

接口：ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

效果: MemoryEffects::Effect{}

结果：

结果	描述
`result`	32位无符号整数

`tti.experimental_fpsan_embed` (triton::instrument::ExperimentalFPSanEmbedOp)

在 fpsan 整数环中嵌入浮点值

语法

operation ::= `tti.experimental_fpsan_embed` $val attr-dict `:` functional-type(operands, $result)

特征：AlwaysSpeculatableImplTrait, Elementwise, SameOperandsAndResultEncoding, SameOperandsAndResultShape

接口: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

效果: MemoryEffects::Effect{}

操作数:

操作数	描述
`val`	浮点数或浮点值的秩张量

结果:

结果	描述
`result`	整数或整数值的秩张量

`tti.experimental_fpsan_unembed` (triton::instrument::ExperimentalFPSanUnembedOp)

将 fpsan 整数载荷解嵌入为浮点值

语法

operation ::= `tti.experimental_fpsan_unembed` $val attr-dict `:` functional-type(operands, $result)

特征：AlwaysSpeculatableImplTrait, Elementwise, SameOperandsAndResultEncoding, SameOperandsAndResultShape

接口: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

效果: MemoryEffects::Effect{}

操作数：

操作数	描述
`val`	整数或整数值的秩张量

结果：

结果	描述
`result`	浮点数或浮点值的秩张量

`tti.experimental_gsan_atomic_cas` (triton::instrument::ExperimentalGSanAtomicCASOp)

降低 GSan 插装的原子比较交换 (CAS) 操作

语法

operation ::= `tti.experimental_gsan_atomic_cas` $sem `,` $scope `,` $ptr `,` $cmp `,` $val attr-dict `:`
              functional-type(operands, $result)

特征：SameOperandsAndResultEncoding, SameOperandsAndResultShape

属性：

属性	MLIR 类型	描述
`sem`	::mlir::triton::MemSemanticAttr	允许的32位无符号整数情况：1, 2, 3, 4
`scope`	::mlir::triton::MemSyncScopeAttr	允许的32位无符号整数情况：1, 2, 3

操作数：

操作数	描述
`ptr`	指针或指针值的秩张量
`cmp`	浮点数或浮点值的秩张量，或整数或整数值的秩张量，或指针或指针值的秩张量
`val`	浮点数或浮点值的秩张量，或整数或整数值的秩张量，或指针或指针值的秩张量

结果：

结果	描述
`result`	浮点数或浮点值的秩张量，或整数或整数值的秩张量，或指针或指针值的秩张量

`tti.experimental_gsan_atomic_rmw` (triton::instrument::ExperimentalGSanAtomicRMWOp)

降低 GSan 插装的原子读-改-写 (RMW) 操作

语法

operation ::= `tti.experimental_gsan_atomic_rmw` $atomic_rmw_op `,` $sem `,` $scope `,` $ptr `,` $val (`,` $mask^)? attr-dict `:`
              functional-type(operands, $result)

特征：SameOperandsAndResultEncoding, SameOperandsAndResultShape

属性：

属性	MLIR 类型	描述
`atomic_rmw_op`	::mlir::triton::RMWOpAttr	允许的 32 位无符号整数值：1, 2, 3, 4, 5, 6, 7, 8, 9, 10
`sem`	::mlir::triton::MemSemanticAttr	允许的32位无符号整数情况：1, 2, 3, 4
`scope`	::mlir::triton::MemSyncScopeAttr	允许的32位无符号整数情况：1, 2, 3

操作数：

操作数	描述
`ptr`	指针或指针值的秩张量
`val`	浮点数或浮点值的秩张量，或整数或整数值的秩张量，或指针或指针值的秩张量
`mask`	1 位无符号整数或 1 位无符号整数值的秩张量

结果：

结果	描述
`result`	浮点数或浮点值的秩张量，或整数或整数值的秩张量，或指针或指针值的秩张量

`tti.experimental_gsan_atomic_tensor_access` (triton::instrument::ExperimentalGSanAtomicTensorAccessOp)

为 GSan 插装张量原子访问

语法

operation ::= `tti.experimental_gsan_atomic_tensor_access` $sem `,` $scope `,` $ptr (`,` $mask^)? attr-dict `:` type($ptr)

为张量指针访问发出运行时插装，其单个元素为原子读-改-写操作。

属性：

属性	MLIR 类型	描述
`sem`	::mlir::triton::MemSemanticAttr	允许的32位无符号整数情况：1, 2, 3, 4
`scope`	::mlir::triton::MemSyncScopeAttr	允许的32位无符号整数情况：1, 2, 3

操作数：

操作数	描述
`ptr`	指针或指针值的秩张量
`mask`	1 位无符号整数或 1 位无符号整数值的秩张量

`tti.experimental_gsan_init` (triton::instrument::ExperimentalGSanInitOp)

初始化 GSan 线程

`tti.experimental_gsan_tensor_access` (triton::instrument::ExperimentalGSanTensorAccessOp)

为 GSan 插装张量加载/存储访问

语法

operation ::= `tti.experimental_gsan_tensor_access` $ptr `,` $isStore (`,` $mask^)? attr-dict `:` type($ptr)

为张量指针访问发出运行时插装。指针和可选掩码由 GSan 运行时消耗。

属性：

属性	MLIR 类型	描述
`isStore`	::mlir::BoolAttr	布尔属性

操作数：

操作数	描述
`ptr`	指针或指针值的秩张量
`mask`	1 位无符号整数或 1 位无符号整数值的秩张量

`tti.experimental_gsan_tensordesc_info` (triton::instrument::ExperimentalGSanTensorDescInfoOp)

从原生张量描述符解码 GSan 描述符元数据

语法

operation ::= `tti.experimental_gsan_tensordesc_info` $desc attr-dict `:` qualified(type($desc)) `->` type($result)

将原生张量描述符解码为底层的基指针、形状和步长值。

特性: AlwaysSpeculatableImplTrait

接口: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

效果: MemoryEffects::Effect{}

操作数：

操作数	描述
`desc`	张量描述符类型

结果：

结果	描述
`result`	任意类型的可变参数

`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}

操作数：

操作数	描述
`锁`	指针或指针值的秩张量
`pred`	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}

操作数:

操作数	描述
`锁`	指针或指针值的秩张量
`pred`	1位无符号整数

`tti.experimental_memdesc_to_i32` (triton::instrument::ExperimentalMemDescToI32Op)

将 memdesc 转换为其 i32 基指针

语法

operation ::= `tti.experimental_memdesc_to_i32` $memdesc attr-dict `:` type($memdesc)

从给定的 memdesc 中提取基指针，并将其作为 32 位整数返回。这可用于根据并发清理器（concurrency sanitizer）维护的屏障指针张量来比较 memdesc。

特性: AlwaysSpeculatableImplTrait

接口：ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

效果: MemoryEffects::Effect{}

操作数：

操作数	描述
`memdesc`	Triton IR 类型系统中的内存描述符类型 (`::mlir::triton::gpu::MemDescType`)

结果：

结果	描述
`result`	32位无符号整数

TritonInstrumentOps

tti.experimental_assert_uniform (triton::instrument::ExperimentalAssertUniformOp)

属性：

操作数：

tti.experimental_buffer_descriptors (triton::instrument::ExperimentalBufferDescriptorsOp)

属性：

结果：

tti.experimental_cluster_cta_id (triton::instrument::ExperimentalClusterCTAIdOp)

结果：

tti.experimental_fpsan_embed (triton::instrument::ExperimentalFPSanEmbedOp)

操作数:

结果:

tti.experimental_fpsan_unembed (triton::instrument::ExperimentalFPSanUnembedOp)

操作数：

结果：

tti.experimental_gsan_atomic_cas (triton::instrument::ExperimentalGSanAtomicCASOp)

属性：

操作数：

结果：

tti.experimental_gsan_atomic_rmw (triton::instrument::ExperimentalGSanAtomicRMWOp)

属性：

操作数：

结果：

tti.experimental_gsan_atomic_tensor_access (triton::instrument::ExperimentalGSanAtomicTensorAccessOp)

属性：

操作数：

tti.experimental_gsan_init (triton::instrument::ExperimentalGSanInitOp)

tti.experimental_gsan_tensor_access (triton::instrument::ExperimentalGSanTensorAccessOp)

属性：

操作数：

tti.experimental_gsan_tensordesc_info (triton::instrument::ExperimentalGSanTensorDescInfoOp)

操作数：

结果：

tti.experimental_lock_acquire (triton::instrument::ExperimentalLockAcquireOp)

操作数：

tti.experimental_lock_release (triton::instrument::ExperimentalLockReleaseOp)

操作数:

tti.experimental_memdesc_to_i32 (triton::instrument::ExperimentalMemDescToI32Op)

操作数：

结果：

`tti.experimental_assert_uniform` (triton::instrument::ExperimentalAssertUniformOp)

`tti.experimental_buffer_descriptors` (triton::instrument::ExperimentalBufferDescriptorsOp)

`tti.experimental_cluster_cta_id` (triton::instrument::ExperimentalClusterCTAIdOp)

`tti.experimental_fpsan_embed` (triton::instrument::ExperimentalFPSanEmbedOp)

`tti.experimental_fpsan_unembed` (triton::instrument::ExperimentalFPSanUnembedOp)

`tti.experimental_gsan_atomic_cas` (triton::instrument::ExperimentalGSanAtomicCASOp)

`tti.experimental_gsan_atomic_rmw` (triton::instrument::ExperimentalGSanAtomicRMWOp)

`tti.experimental_gsan_atomic_tensor_access` (triton::instrument::ExperimentalGSanAtomicTensorAccessOp)

`tti.experimental_gsan_init` (triton::instrument::ExperimentalGSanInitOp)

`tti.experimental_gsan_tensor_access` (triton::instrument::ExperimentalGSanTensorAccessOp)

`tti.experimental_gsan_tensordesc_info` (triton::instrument::ExperimentalGSanTensorDescInfoOp)

`tti.experimental_lock_acquire` (triton::instrument::ExperimentalLockAcquireOp)

`tti.experimental_lock_release` (triton::instrument::ExperimentalLockReleaseOp)

`tti.experimental_memdesc_to_i32` (triton::instrument::ExperimentalMemDescToI32Op)