TritonGPUOps

`ttg.async_commit_group` (triton::gpu::AsyncCommitGroupOp)

将待处理的异步拷贝提交到一个可等待的异步组中

语法

operation ::= `ttg.async_commit_group` (`tokens` $inputTokens^)? attr-dict

关闭当前的 async_copy_* 操作批次，并允许通过 ttg.async_wait 对其进行等待。这是确保异步拷贝操作可被等待所必需的。

特征 (Traits): VerifyMemDescLayoutsTrait, VerifyTensorLayoutsTrait

接口 (Interfaces): InferTypeOpInterface

操作数：

操作数	描述
`输入标记 (inputTokens)`	异步标记类型的可变参数

结果：

结果	描述
`异步标记 (asyncToken)`	异步令牌类型

`ttg.async_copy_global_to_local` (triton::gpu::AsyncCopyGlobalToLocalOp)

异步将数据从全局内存拷贝到局部内存

语法

operation ::= `ttg.async_copy_global_to_local` $src `,` $result (`mask` $mask^)? (`other` $other^)?
              oilist(`cacheModifier` `=` $cache | `evictionPolicy` `=` $evict)
              attr-dict `:` type($src) `->` type($result)

该操作异步将数据从全局内存拷贝到局部内存。它类似于 tt.load，区别在于数据被拷贝到由内存描述符指向的局部内存，而不是分布式张量。其余操作数与 tt.load 相同。连续性 (Contiguity) 是在给定布局和掩码下，单次向量化操作中可以加载的最大元素数量。即使无法根据 IR 证明对齐方式，这也允许操作使用 async_copy_global_to_local。

只有在发出 ttg.async_wait 以等待 async_copy_global_to_local 完成后，数据才会在局部内存中可用。必须使用 ttg.async_commit_group 提交异步拷贝操作以关闭批次，从而允许对其进行等待。

特征: AttrSizedOperandSegments, VerifyMemDescLayoutsTrait, VerifyTensorLayoutsTrait

接口: InferTypeOpInterface, PredicatedOpInterface

属性：

属性	MLIR 类型	描述
`cache`	::mlir::triton::CacheModifierAttr	允许的32位无符号整数情况：1, 2, 3, 4, 5, 6, 7
`evict`	::mlir::triton::EvictionPolicyAttr	允许的32位无符号整数情况：1, 2, 3
`isVolatile`	::mlir::BoolAttr	布尔属性
`连续性 (contiguity)`	::mlir::IntegerAttr	32位无符号整数属性

操作数:

操作数	描述
`src`	指针值的秩张量
`result`	Triton IR 类型系统中的内存描述符类型 (`::mlir::triton::gpu::MemDescType`)
`mask`	1位无符号整数值的张量
`other`	浮点数或浮点值的秩张量，或整数或整数值的秩张量，或指针或指针值的秩张量

结果：

结果	描述
`token`	异步令牌类型

`ttg.async_wait` (triton::gpu::AsyncWaitOp)

确保所有指定的 async_copy* 操作均已完成。

语法

operation ::= `ttg.async_wait` ($asyncToken^)? attr-dict

async_wait 操作会等待，直到最多有“num”个异步拷贝组处于待处理状态，且不会同步 CTA 执行。它接收零个或多个 asyncToken 以及一个整数 num，用于指定在 async_wait 操作完成后允许有多少个异步拷贝组保持待处理状态。num = 0 表示等待直到所有异步拷贝组完成。

此操作不提供 CTA 中的任何同步；如果需要同步，请在此操作之外使用 ttg.local_barrier。

特征: MemWaitOpTrait, VerifyMemDescLayoutsTrait, VerifyTensorLayoutsTrait

接口 (Interfaces): InferTypeOpInterface

属性：

属性	MLIR 类型	描述
`数量 (num)`	::mlir::IntegerAttr	32位无符号整数属性

操作数：

操作数	描述
`异步标记 (asyncToken)`	异步标记类型的可变参数

结果:

结果	描述
`返回标记 (retToken)`	异步令牌类型

`ttg.barrier` (triton::gpu::BarrierOp)

同步 CTA 中所有线程对所选地址空间的执行和读/写操作。

barrier 操作同步 CTA 中所有线程在所选地址空间之间的执行和所有操作。它用于协调 CTA 中线程之间的通信。

此操作会等待直到 CTA 中的所有线程都到达 barrier（用于同步），并且这些线程在操作之前对所选地址空间所做的操作对 CTA 中的所有线程可见。

通过在这些访问之间使用 barrier 同步指定的范围，可以避免访问同一内存的线程之间的数据冲突。

barrier 操作仅在 CTA 的选定地址空间上提供同步和内存保证。

强制性的 addrspace 属性是一个位掩码，描述了在 barrier 完成时哪些地址空间将是可见的。

none 仅控制同步（无内存排序）。
local 共享内存操作已完成并对整个 CTA 可见。
global_read 全局内存读取操作已完成并对整个 CTA 可见。
global_write 全局内存写入操作已完成并对整个 CTA 可见。
tensor_read 张量内存读取操作已完成并对整个 CTA 可见。
tensor_write 张量内存写入操作已完成并对整个 CTA 可见。
all ["local", "global_read", "global_write", "tensor_read", "tensor_write"] 的便捷别名。

多个地址空间可以组合（例如 local|tensor_write）。none 不能与其他地址空间组合。

示例

ttg.barrier local
ttg.barrier local|global_read|global_write

特征 (Traits): VerifyMemDescLayoutsTrait, VerifyTensorLayoutsTrait

属性：

属性	MLIR 类型	描述
`地址空间 (addrSpace)`	::mlir::triton::gpu::AddrSpaceAttr

`ttg.convert_layout` (triton::gpu::ConvertLayoutOp)

转换布局

语法

operation ::= `ttg.convert_layout` $src attr-dict `:` type($src) `->` type($result)

特征: AlwaysSpeculatableImplTrait, SameOperandsAndResultElementType, SameOperandsAndResultShape, VerifyMemDescLayoutsTrait, VerifyTensorLayoutsTrait

接口: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

效果: MemoryEffects::Effect{}

操作数：

操作数	描述
`src`	浮点数、整数或 ptr 值的秩张量

结果：

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

`ttg.fp4_to_fp` (triton::gpu::Fp4ToFpOp)

将 fp4 (e2m1) 转换为 fp

语法

operation ::= `ttg.fp4_to_fp` $src attr-dict `:` type($src) `->` type($result)

将表示为打包 i8 的 fp4 (e2m1) 转换为 fp。

i8 的低 4 位表示第一个 fp4 元素，高 4 位表示第二个 fp4 元素。

axis 属性指定了 fp4 元素打包所沿的轴。

特征: AlwaysSpeculatableImplTrait, VerifyMemDescLayoutsTrait, VerifyTensorLayoutsTrait

接口: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

效果: MemoryEffects::Effect{}

属性:

属性	MLIR 类型	描述
`axis`	::mlir::IntegerAttr	32位无符号整数属性

操作数：

操作数	描述
`src`	8位无符号整数值的秩张量

结果：

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

`ttg.global_scratch_alloc` (triton::gpu::GlobalScratchAllocOp)

分配全局内存缓冲区

语法

operation ::= `ttg.global_scratch_alloc` attr-dict `:` qualified(type($result))

此操作在全局内存中分配一个对当前程序私有的缓冲区。可以使用可选的 third_party_allocation 单位属性标记自定义的第三方分配。持有集群范围状态的并发清理分配可以使用可选的 shared_cluster_state 单位属性进行标记。

特征 (Traits): VerifyMemDescLayoutsTrait, VerifyTensorLayoutsTrait

属性：

属性	MLIR 类型	描述
`字节数 (nbytes)`	::mlir::IntegerAttr	32位无符号整数属性
`对齐方式 (alignment)`	::mlir::IntegerAttr	32位无符号整数属性
`第三方分配 (third_party_allocation)`	::mlir::UnitAttr	单元属性
`共享集群状态 (shared_cluster_state)`	::mlir::UnitAttr	单元属性

结果:

结果	描述
`result`	ptr

`ttg.local_alloc` (triton::gpu::LocalAllocOp)

分配张量

语法

operation ::= `ttg.local_alloc` ($src^)? attr-dict `:` functional-type(operands, results)

此操作在共享内存中分配缓冲区，并返回一个包含地址和缓冲区视图的描述符。

显式释放缓冲区是可选的；请参阅 local_dealloc。

src 操作数是所分配缓冲区的可选初始化器。它必须与缓冲区具有相同的元素类型。如果未指定 src，则返回的缓冲区必须是可变的。

特征 (Traits): VerifyMemDescLayoutsTrait, VerifyTensorLayoutsTrait

接口： MemoryEffectOpInterface

属性：

属性	MLIR 类型	描述
`对齐方式 (alignment)`	::mlir::IntegerAttr	32位无符号整数属性

操作数：

操作数	描述
`src`	浮点数、整数或 ptr 值的秩张量

结果：

结果	描述
`result`	Triton IR 类型系统中的内存描述符类型 (`::mlir::triton::gpu::MemDescType`)

`ttg.local_atomic_scatter_rmw` (triton::gpu::LocalAtomicScatterRMWOp)

沿指定轴原子地将 RMW 元素散布到共享内存中

语法

operation ::= `ttg.local_atomic_scatter_rmw` $atomic_rmw_op `,` $dst `[` $indices `]` `,` $values (`,` $mask^)? attr-dict `:`
              functional-type(operands, results)

使用沿单个指定轴的索引张量，原子地更新共享内存描述符中的元素。值张量与索引张量具有相同的形状，结果返回在每个更新位置观察到的先前值。

对于每个输入位置 I，该操作原子地更新 dst，其中轴上的坐标被 indices[I] 替换：dst[I[0], …, indices[I], …, I[n]] = rmw(dst[I[0], …, indices[I], …, I[n]], values[I])

如果存在掩码，则仅对 mask[I] 为真的位置执行更新。被屏蔽的位置不会更新共享内存，并返回未定义的旧值。

特征 (Traits): VerifyMemDescLayoutsTrait, VerifyTensorLayoutsTrait

接口 (Interfaces): InferTypeOpInterface

属性：

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

操作数：

操作数	描述
`dst`	Triton IR 类型系统中的内存描述符类型 (`::mlir::triton::gpu::MemDescType`)
`值 (values)`	浮点数、整数或 ptr 值的秩张量
`indices`	整数值的秩张量
`mask`	1位无符号整数值的张量

结果：

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

`ttg.local_dealloc` (triton::gpu::LocalDeallocOp)

释放缓冲区

语法

operation ::= `ttg.local_dealloc` $src attr-dict `:` qualified(type($src))

此操作显式释放缓冲区。在此操作之后使用该缓冲区是未定义的行为。

此操作是可选的。如果您没有显式释放缓冲区，编译器会假定它在后支配所有分配使用的第一个点处被释放。

因为我们假定内存描述符在后支配其使用的第一个点处失效，所以等待内存描述符上异步操作完成的操作（例如 ttng.warp_group_dot_wait）也应将内存描述符作为操作数。

特征 (Traits): VerifyMemDescLayoutsTrait, VerifyTensorLayoutsTrait

操作数:

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

`ttg.local_gather` (triton::gpu::LocalGatherOp)

沿指定轴从共享内存中聚集元素

语法

operation ::= `ttg.local_gather` $src `[` $indices `]` (`token` $token^)? attr-dict `:` qualified(type($src)) `,` type($indices) `->` type($result)

使用沿单个指定轴的索引张量从共享内存描述符中聚集元素。输出张量与索引张量具有相同的形状。

对于每个输出位置 I，该操作从 src 读取数据，其中聚集轴上的坐标被 indices[I] 替换：result[I] = src[I[0], …, indices[I], …, I[n]]，其中轴维度被索引值替换。

这与 tt.gather 的行为匹配，但在共享内存描述符上运行。

特征: LocalLoadTrait, VerifyMemDescLayoutsTrait, VerifyTensorLayoutsTrait

属性:

属性	MLIR 类型	描述
`axis`	::mlir::IntegerAttr	32位无符号整数属性

操作数:

操作数	描述
`src`	Triton IR 类型系统中的内存描述符类型 (`::mlir::triton::gpu::MemDescType`)
`indices`	整数值的秩张量
`token`	异步令牌类型

结果:

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

`ttg.local_load` (triton::gpu::LocalLoadOp)

将缓冲区从局部内存加载到分布式张量中

语法

operation ::= `ttg.local_load` $src (`token` $token^)? attr-dict `:` qualified(type($src)) `->` type($result)

将张量从局部内存描述符加载到分布式张量中。

特征: LocalLoadTrait, VerifyMemDescLayoutsTrait, VerifyTensorLayoutsTrait

操作数:

操作数	描述
`src`	Triton IR 类型系统中的内存描述符类型 (`::mlir::triton::gpu::MemDescType`)
`token`	异步令牌类型

结果:

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

`ttg.local_scatter` (triton::gpu::LocalScatterOp)

沿指定轴将元素散布到共享内存中

语法

operation ::= `ttg.local_scatter` $dst `[` $indices `]` `,` $values (`token` $token^)? attr-dict `:` qualified(type($dst)) `,` type($indices) `,` type($values)

使用沿单个指定轴的索引张量将元素散布到共享内存描述符。值张量与索引张量具有相同的形状。

对于每个输入位置 I，该操作写入 dst，其中散布轴上的坐标被 indices[I] 替换：dst[I[0], …, indices[I], …, I[n]] = values[I]，其中轴维度被索引值替换。

这是 local_gather 的逆操作，并在运行时计算出的索引处写入共享内存。

特征 (Traits): VerifyMemDescLayoutsTrait, VerifyTensorLayoutsTrait

属性:

属性	MLIR 类型	描述
`axis`	::mlir::IntegerAttr	32位无符号整数属性

操作数：

操作数	描述
`dst`	Triton IR 类型系统中的内存描述符类型 (`::mlir::triton::gpu::MemDescType`)
`值 (values)`	浮点数、整数或 ptr 值的秩张量
`indices`	整数值的秩张量
`token`	异步令牌类型

`ttg.local_store` (triton::gpu::LocalStoreOp)

将分布式张量存储到局部内存的缓冲区中

语法

operation ::= `ttg.local_store` $src `,` $dst attr-dict `:` type($src) `->` qualified(type($dst))

将分布式张量存储到局部内存的缓冲区中。

特征 (Traits): VerifyMemDescLayoutsTrait, VerifyTensorLayoutsTrait

操作数:

操作数	描述
`src`	浮点数、整数或 ptr 值的秩张量
`dst`	Triton IR 类型系统中的内存描述符类型 (`::mlir::triton::gpu::MemDescType`)

`ttg.mask` (triton::gpu::MaskOp)

用于流水线化的掩码操作

特征: SingleBlock, VerifyMemDescLayoutsTrait, VerifyTensorLayoutsTrait

操作数:

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

结果:

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

`ttg.mask.return` (triton::gpu::MaskReturnOp)

掩码运算符的终止符

语法

operation ::= `ttg.mask.return` $result attr-dict `:` type($result)

特征: AlwaysSpeculatableImplTrait, HasParent<MaskOp>, ReturnLike, Terminator, VerifyMemDescLayoutsTrait, VerifyTensorLayoutsTrait

接口 (Interfaces): ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface), RegionBranchTerminatorOpInterface

效果: MemoryEffects::Effect{}

操作数:

操作数	描述
`result`	任意类型的可变参数

`ttg.memdesc_index` (triton::gpu::MemDescIndexOp)

获取描述符的子视图。

语法

operation ::= `ttg.memdesc_index` $src `[` $index `]` attr-dict `:` qualified(type($src)) `->` qualified(type($result))

此操作返回一个指向输入描述符第 0 维第 i 个元素的新描述符。
它不会影响底层内存。
例如，假设

输入形状为 2x4x16xf16，

输出形状为 4x16xf16，并且

索引 = 1。那么输出描述符等同于 input[1]，其中 input 是逻辑张量。

特征: AlwaysSpeculatableImplTrait, MemDescViewTrait, VerifyMemDescLayoutsTrait, VerifyTensorLayoutsTrait

接口: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

效果: MemoryEffects::Effect{}

操作数:

操作数

描述

src

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

索引 (index)

32位无符号整数

结果:

结果

描述

result

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

ttg.memdesc_reinterpret (triton::gpu::MemDescReinterpretOp)

将内存描述符重解释为不同的类型和形状

语法

operation ::= `ttg.memdesc_reinterpret` $src attr-dict `:` qualified(type($src)) `->` qualified(type($result))

ttg.memdesc_reinterpret 操作将内存描述符重解释为具有不同形状和元素类型的描述符。由于内存描述符缺乏步长 (strides)，此操作仅在原始内存描述符是连续的情况下有效。不允许对子视图进行重解释；请先重解释父描述符，然后再获取重解释后描述符的子视图。
特征: AlwaysSpeculatableImplTrait, MemDescViewTrait, VerifyMemDescLayoutsTrait, VerifyTensorLayoutsTrait

接口: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

效果: MemoryEffects::Effect{}

操作数:

操作数

描述

src

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

结果:

结果

描述

result

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

ttg.memdesc_reshape (triton::gpu::MemDescReshapeOp)

为新形状创建描述符

语法

operation ::= `ttg.memdesc_reshape` $src attr-dict `:` qualified(type($src)) `->` qualified(type($result))

此操作返回一个代表底层缓冲区重塑视图的新描述符。这不会影响内存。
特征: AlwaysSpeculatableImplTrait, MemDescViewTrait, SameOperandsAndResultElementType, VerifyMemDescLayoutsTrait, VerifyTensorLayoutsTrait

接口: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

效果: MemoryEffects::Effect{}

操作数：

操作数

描述

src

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

结果:

结果

描述

result

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

ttg.memdesc_subslice (triton::gpu::MemDescSubsliceOp)

获取描述符的子视图。

语法

operation ::= `ttg.memdesc_subslice` $src `[` custom<Offsets>($offsets) `]` attr-dict `:` qualified(type($src)) `->` qualified(type($result))

此操作返回一个代表逻辑张量子视图的新描述符。它不会影响底层内存。
例如，假设

输入形状为 32x16xf16，

输出形状为 8x16xf16，并且

偏移量 = [2, 1]。那么在 Python 语法中，子视图覆盖 input[2:8+2, 1:16+1]，其中 input 是逻辑张量。

偏移量必须大于或等于张量的平铺大小（或为零）。
特征: AlwaysSpeculatableImplTrait, MemDescViewTrait, VerifyMemDescLayoutsTrait, VerifyTensorLayoutsTrait

接口: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

效果: MemoryEffects::Effect{}

属性:

属性 MLIR 类型描述

offsets ::mlir::DenseI32ArrayAttr i32 密集数组属性

操作数：

操作数

描述

src

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

结果:

结果

描述

result

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

ttg.memdesc_trans (triton::gpu::MemDescTransOp)

转置描述符

语法

operation ::= `ttg.memdesc_trans` $src attr-dict `:` qualified(type($src)) `->` qualified(type($result))

此操作返回一个代表缓冲区转置视图的新描述符。
特征: AlwaysSpeculatableImplTrait, InferTypeOpAdaptor, MemDescViewTrait, SameOperandsAndResultElementType, VerifyMemDescLayoutsTrait, VerifyTensorLayoutsTrait

接口: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface), TransposeOpInterface

效果: MemoryEffects::Effect{}

属性:

属性 MLIR 类型描述

顺序 (order) ::mlir::DenseI32ArrayAttr i32 密集数组属性

操作数:

操作数

描述

src

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

结果:

结果

描述

result

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

ttg.predicate_stage (triton::gpu::PredicateStageOp)

流水线阶段谓词

语法

operation ::= `ttg.predicate_stage` $iv `,` $ub `,` $step `maxStage` $maxStage `stage` $stage attr-dict `:` type($iv) `->` type($result)

特征: AlwaysSpeculatableImplTrait, VerifyMemDescLayoutsTrait, VerifyTensorLayoutsTrait

接口：ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

效果: MemoryEffects::Effect{}

属性:

属性 MLIR 类型描述

最大阶段 (maxStage) ::mlir::IntegerAttr 32位无符号整数属性

阶段 (stage) ::mlir::IntegerAttr 32位无符号整数属性

操作数:

操作数

描述

迭代变量 (iv)

无符号整数或索引

上限 (ub)

无符号整数或索引

步长 (step)

无符号整数或索引

结果:

结果

描述

result

1位无符号整数

ttg.warp_id (triton::gpu::WarpIdOp)

返回 GPU 线程束 ID (Warp ID)

语法

operation ::= `ttg.warp_id` attr-dict

此操作返回 GPU 线程束 ID。如果存在硬件寄存器，这可以转换为读取硬件寄存器，或者仅将线程 ID 除以线程束大小。
omitUniformHint 属性指示在 NVIDIA 后端中是否省略为 LLVM 发出 nvvm.shfl.sync idx 0。
特征: AlwaysSpeculatableImplTrait, VerifyMemDescLayoutsTrait, VerifyTensorLayoutsTrait

接口：ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

效果: MemoryEffects::Effect{}

属性：

属性 MLIR 类型描述

省略统一提示 (omitUniformHint) ::mlir::UnitAttr 单元属性

结果:

结果

描述

result

32位无符号整数

ttg.warp_return (triton::gpu::WarpReturnOp)

分区区域的隐式终止符

语法

operation ::= `ttg.warp_return` attr-dict

ttg.warp_return 操作是结束 ttg.warp_specialize 操作分区区域的隐式终止符。它没有操作数，因为这些区域不能返回任何内容。
待办事项：支持从分区区域返回统一值。
特征: AlwaysSpeculatableImplTrait, HasParent<WarpSpecializePartitionsOp>, ReturnLike, Terminator, VerifyMemDescLayoutsTrait, VerifyTensorLayoutsTrait
接口 (Interfaces): ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface), RegionBranchTerminatorOpInterface

效果: MemoryEffects::Effect{}

ttg.warp_specialize (triton::gpu::WarpSpecializeOp)

在多个线程束组 (warpgroups) 上异步执行代码

ttg.warp_specialize 操作代表在不同的线程束组上同时执行不同的代码。线程束组是 2 的幂次线程束的集合，其线程束数量可能与包含区域中的不同。
该操作的“默认”区域代表当前执行的线程束组所执行的代码。此区域允许隐式捕获。该操作包含多个与上方隔离的“分区”区域。它们必须是隔离的，因为这些区域代表不同的布局域，因为线程束的数量不同。
从语义上讲，每个区域的执行对于每个线程束组都是同时开始的，并且所有线程束组在操作结束时汇合。
示例

%0 = ttg.warp_specialize(%a, %b) default { %out = some_operation(%a) // implicit capture of `%a` ttg.warp_yield %out : i32 } partition0(%arg0: i32, %arg1: i32) num_warps(8) { some_async_dispatch(%arg0, %arg1) ttg.warp_return } partition1(%arg0: i32, %arg1: i32) num_warps(1) { some_async_dispatch(%arg0, %arg1) ttg.warp_return } : (i32, i32) -> i32

特征: AsyncRegions, RecursiveMemoryEffects, RecursivelySpeculatableImplTrait, VerifyMemDescLayoutsTrait, VerifyTensorLayoutsTrait
接口: ConditionallySpeculatable, RegionBranchOpInterface

属性:

属性 MLIR 类型描述

分区线程束数量 (partitionNumWarps) ::mlir::DenseI32ArrayAttr i32 密集数组属性

线程束组起始 ID (warpGroupStartIds) ::mlir::DenseI32ArrayAttr i32 密集数组属性

请求的寄存器 (requestedRegisters) ::mlir::DenseI32ArrayAttr i32 密集数组属性

实际寄存器 (actualRegisters) ::mlir::DenseI32ArrayAttr i32 密集数组属性

结果:

结果

描述

默认透传 (defaultPassthrough)

任意类型的可变参数

ttg.warp_specialize.partitions (triton::gpu::WarpSpecializePartitionsOp)

ttg.warp_specialize 的容器操作

由于 MLIR 要求整个操作必须与上方隔离，此操作包含了 ttg.warp_specialize 的实际隔离区域。
特征: HasParent<WarpSpecializeOp>, IsolatedFromAbove, RecursiveMemoryEffects, RecursivelySpeculatableImplTrait, Terminator, VerifyMemDescLayoutsTrait, VerifyTensorLayoutsTrait
接口: ConditionallySpeculatable, RegionBranchOpInterface

操作数:

操作数

描述

显式捕获 (explicitCaptures)

任意类型的可变参数

ttg.warp_yield (triton::gpu::WarpYieldOp)

从 ttg.warp_specialize 的默认区域中产生 (Yield)

语法

operation ::= `ttg.warp_yield` ($values^)? attr-dict (`:` type($values)^)?

ttg.warp_yield 操作是 ttg.warp_specialize 操作“默认”区域的终止符。操作数被透明地传递为 ttg.warp_specialize 操作的 SSA 结果。
示例

ttg.warp_yield %a, %b : i32, tensor<32xbf16, #blocked>

特征: AlwaysSpeculatableImplTrait, HasParent<WarpSpecializeOp>, ReturnLike, Terminator, VerifyMemDescLayoutsTrait, VerifyTensorLayoutsTrait
接口 (Interfaces): ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface), RegionBranchTerminatorOpInterface

效果: MemoryEffects::Effect{}

操作数:

操作数

描述

值 (values)

任意类型的可变参数

TritonGPUOps

ttg.async_commit_group (triton::gpu::AsyncCommitGroupOp)

操作数：

结果：

ttg.async_copy_global_to_local (triton::gpu::AsyncCopyGlobalToLocalOp)

属性：

操作数:

结果：

ttg.async_wait (triton::gpu::AsyncWaitOp)

属性：

操作数：

结果:

ttg.barrier (triton::gpu::BarrierOp)

属性：

ttg.convert_layout (triton::gpu::ConvertLayoutOp)

操作数：

结果：

ttg.fp4_to_fp (triton::gpu::Fp4ToFpOp)

属性:

操作数：

结果：

ttg.global_scratch_alloc (triton::gpu::GlobalScratchAllocOp)

属性：

结果:

ttg.local_alloc (triton::gpu::LocalAllocOp)

属性：

操作数：

结果：

ttg.local_atomic_scatter_rmw (triton::gpu::LocalAtomicScatterRMWOp)

属性：

操作数：

结果：

ttg.local_dealloc (triton::gpu::LocalDeallocOp)

操作数:

ttg.local_gather (triton::gpu::LocalGatherOp)

属性:

操作数:

结果:

ttg.local_load (triton::gpu::LocalLoadOp)

操作数:

结果:

ttg.local_scatter (triton::gpu::LocalScatterOp)

属性:

操作数：

ttg.local_store (triton::gpu::LocalStoreOp)

操作数:

ttg.mask (triton::gpu::MaskOp)

操作数:

结果:

ttg.mask.return (triton::gpu::MaskReturnOp)

操作数:

ttg.memdesc_index (triton::gpu::MemDescIndexOp)

操作数:

结果:

ttg.memdesc_reinterpret (triton::gpu::MemDescReinterpretOp)

操作数:

结果:

ttg.memdesc_reshape (triton::gpu::MemDescReshapeOp)

操作数：

结果:

ttg.memdesc_subslice (triton::gpu::MemDescSubsliceOp)

属性:

操作数：

结果:

ttg.memdesc_trans (triton::gpu::MemDescTransOp)

属性:

操作数:

结果:

ttg.predicate_stage (triton::gpu::PredicateStageOp)

属性:

操作数:

结果:

ttg.warp_id (triton::gpu::WarpIdOp)

属性：

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

ttg.warp_return (triton::gpu::WarpReturnOp)

ttg.warp_specialize (triton::gpu::WarpSpecializeOp)

属性:

结果: 结果 描述 默认透传 (defaultPassthrough) 任意类型的可变参数

ttg.warp_specialize.partitions (triton::gpu::WarpSpecializePartitionsOp)

`ttg.async_commit_group` (triton::gpu::AsyncCommitGroupOp)

`ttg.async_copy_global_to_local` (triton::gpu::AsyncCopyGlobalToLocalOp)

`ttg.async_wait` (triton::gpu::AsyncWaitOp)

`ttg.barrier` (triton::gpu::BarrierOp)

`ttg.convert_layout` (triton::gpu::ConvertLayoutOp)

`ttg.fp4_to_fp` (triton::gpu::Fp4ToFpOp)

`ttg.global_scratch_alloc` (triton::gpu::GlobalScratchAllocOp)

`ttg.local_alloc` (triton::gpu::LocalAllocOp)

`ttg.local_atomic_scatter_rmw` (triton::gpu::LocalAtomicScatterRMWOp)

`ttg.local_dealloc` (triton::gpu::LocalDeallocOp)

`ttg.local_gather` (triton::gpu::LocalGatherOp)

`ttg.local_load` (triton::gpu::LocalLoadOp)

`ttg.local_scatter` (triton::gpu::LocalScatterOp)

`ttg.local_store` (triton::gpu::LocalStoreOp)

`ttg.mask` (triton::gpu::MaskOp)

`ttg.mask.return` (triton::gpu::MaskReturnOp)

`ttg.memdesc_index` (triton::gpu::MemDescIndexOp)

`ttg.memdesc_reinterpret` (triton::gpu::MemDescReinterpretOp)

`ttg.memdesc_reshape` (triton::gpu::MemDescReshapeOp)

`ttg.memdesc_subslice` (triton::gpu::MemDescSubsliceOp)

`ttg.memdesc_trans` (triton::gpu::MemDescTransOp)

`ttg.predicate_stage` (triton::gpu::PredicateStageOp)

`ttg.warp_id` (triton::gpu::WarpIdOp)

结果:

结果

描述

`result`

32位无符号整数

`ttg.warp_return` (triton::gpu::WarpReturnOp)

`ttg.warp_specialize` (triton::gpu::WarpSpecializeOp)

结果:

结果

描述

`默认透传 (defaultPassthrough)`

任意类型的可变参数

`ttg.warp_specialize.partitions` (triton::gpu::WarpSpecializePartitionsOp)

操作数:

操作数

描述

`显式捕获 (explicitCaptures)`

任意类型的可变参数

`ttg.warp_yield` (triton::gpu::WarpYieldOp)

操作数:

操作数

描述

`值 (values)`

任意类型的可变参数