TritonNvidiaGPUOps¶

`ttng.arrive_barrier` (triton::nvidia_gpu::ArriveBarrierOp)¶

对mbarrier执行到达操作

语法：

operation ::= `ttng.arrive_barrier` $alloc `,` $count (`,` $pred^)? attr-dict `:` qualified(type($alloc))

ttng.arrive_barrier 操作在共享内存中的 mbarrier 对象上执行"到达"操作。该操作要求 count 属性至少为1，并将 mbarrier 的待处理到达计数减少指定的计数值。

该操作接受一个可选的谓词。

示例：

ttng.arrive_barrier %barrier, 2 : !ttg.memdesc<1xi64, #shared, #smem, mutable>
ttng.arrive_barrier %barrier, 1, %pred : !ttg.memdesc<1xi64, #shared, #smem, mutable>

特性: VerifyTensorLayoutsTrait

属性:¶

属性	MLIR类型	描述
`count`	::mlir::IntegerAttr	32-bit signless integer attribute

操作数:¶

操作数	描述
`alloc`	Triton IR类型系统中的内存描述符类型 (`::mlir::triton::gpu::MemDescType`)
`pred`	1位无符号整数

`ttng.async_tma_copy_global_to_local` (triton::nvidia_gpu::AsyncTMACopyGlobalToLocalOp)¶

基于描述符从全局内存异步复制数据到本地内存

语法：

operation ::= `ttng.async_tma_copy_global_to_local` $desc_ptr `[` $coord `]` $result `,` $barrier `,` $pred
              oilist(`cacheModifier` `=` $cache | `evictionPolicy` `=` $evict)
              attr-dict `:` qualified(type($desc_ptr)) `,` qualified(type($barrier)) `->` qualified(type($result))

该操作将数据从全局内存异步复制到本地内存。这与tt.load类似，不同之处在于数据被复制到由内存描述符指向的本地内存，而不是分布式张量。复制的数据取决于由desc_ptr指向的全局内存描述符。

特性: VerifyTensorLayoutsTrait

属性:¶

属性	MLIR类型	描述
`cache`	::mlir::triton::CacheModifierAttr	allowed 32-bit signless integer cases: 1, 2, 3, 4, 5, 6, 7
`evict`	::mlir::triton::EvictionPolicyAttr	allowed 32-bit signless integer cases: 1, 2, 3
`isVolatile`	::mlir::BoolAttr	bool attribute

操作数:¶

操作数	描述
`desc_ptr`	Triton IR类型系统中的指针类型 (`::mlir::triton::PointerType`)
`coord`	32位无符号整数的可变参数
`barrier`	Triton IR类型系统中的内存描述符类型 (`::mlir::triton::gpu::MemDescType`)
`result`	Triton IR类型系统中的内存描述符类型 (`::mlir::triton::gpu::MemDescType`)
`pred`	1位无符号整数

`ttng.async_tma_copy_local_to_global` (triton::nvidia_gpu::AsyncTMACopyLocalToGlobalOp)¶

根据描述符将数据从本地内存异步复制到全局内存

语法：

operation ::= `ttng.async_tma_copy_local_to_global` $desc_ptr `[` $coord `]` $src
              attr-dict `:` qualified(type($desc_ptr)) `,` qualified(type($src))

该操作将数据从本地内存异步复制到全局内存。这与tt.store类似，不同之处在于数据是从内存描述符指向的本地内存复制，而不是从分布式张量复制。复制的数据取决于desc_ptr所指向的全局内存描述符。

特性: VerifyTensorLayoutsTrait

操作数:¶

操作数	描述
`desc_ptr`	Triton IR类型系统中的指针类型(`::mlir::triton::PointerType`)
`coord`	32位无符号整数的可变参数
`src`	Triton IR类型系统中的内存描述符类型 (`::mlir::triton::gpu::MemDescType`)

`ttng.async_tma_gather` (triton::nvidia_gpu::AsyncTMAGatherOp)¶

根据描述符从全局内存异步收集数据到本地内存

语法：

operation ::= `ttng.async_tma_gather` $desc_ptr `[` $x_offsets `,` $y_offset `]` $result `,` $barrier `,` $pred
              attr-dict `:` type(operands)

该操作异步地将多行数据从全局内存矩阵收集到本地内存。这与async_tma_copy_global_to_local类似，区别在于每行数据都是独立索引的。

特性: VerifyTensorLayoutsTrait

操作数:¶

操作数	描述
`desc_ptr`	Triton IR类型系统中的指针类型 (`::mlir::triton::PointerType`)
`x_offsets`	32位无符号整数值的排序张量
`y_offset`	32位无符号整数
`barrier`	Triton IR类型系统中的内存描述符类型 (`::mlir::triton::gpu::MemDescType`)
`result`	Triton IR类型系统中的内存描述符类型 (`::mlir::triton::gpu::MemDescType`)
`pred`	1位无符号整数

`ttng.async_tma_scatter` (triton::nvidia_gpu::AsyncTMAScatterOp)¶

基于描述符将本地内存中的数据异步分散到全局内存中

语法：

operation ::= `ttng.async_tma_scatter` $desc_ptr `[` $x_offsets `,` $y_offset `]` $src
              attr-dict `:` type(operands)

ttng.async_tma_scatter 操作将本地内存中多个独立索引的数据行异步分散到全局内存中。该操作将共享内存中的二维张量按照核心张量块nvmma_shared布局分散到全局内存中指定y偏移处的独立索引行。

特性: VerifyTensorLayoutsTrait

操作数:¶

操作数	描述
`desc_ptr`	Triton IR类型系统中的指针类型(`::mlir::triton::PointerType`)
`x_offsets`	32位无符号整数值的排序张量
`y_offset`	32位无符号整数
`src`	Triton IR类型系统中的内存描述符类型 (`::mlir::triton::gpu::MemDescType`)

`ttng.barrier_expect` (triton::nvidia_gpu::BarrierExpectOp)¶

发出一个预期要复制的字节数的屏障信号。

语法：

operation ::= `ttng.barrier_expect` $alloc `,` $size attr-dict `,` $pred `:` qualified(type($alloc))

这个信号屏障表示预期要复制size字节。关联的屏障等待将阻塞，直到预期的字节数被复制完成。

特性: VerifyTensorLayoutsTrait

属性:¶

属性	MLIR类型	描述
`size`	::mlir::IntegerAttr	32-bit signless integer attribute

操作数:¶

操作数	描述
`alloc`	Triton IR类型系统中的内存描述符类型 (`::mlir::triton::gpu::MemDescType`)
`pred`	1位无符号整数

`ttng.cluster_arrive` (triton::nvidia_gpu::ClusterArriveOp)¶

语法：

operation ::= `ttng.cluster_arrive` attr-dict

特性: VerifyTensorLayoutsTrait

属性:¶

属性	MLIR类型	描述
`relaxed`	::mlir::IntegerAttr	1-bit signless integer attribute

`ttng.cluster_wait` (triton::nvidia_gpu::ClusterWaitOp)¶

语法：

operation ::= `ttng.cluster_wait` attr-dict

特性: VerifyTensorLayoutsTrait

`ttng.fence_async_shared` (triton::nvidia_gpu::FenceAsyncSharedOp)¶

异步栅栏代理

语法：

operation ::= `ttng.fence_async_shared` attr-dict

特性: VerifyTensorLayoutsTrait

属性:¶

属性	MLIR类型	描述
`bCluster`	::mlir::BoolAttr	bool attribute

`ttng.init_barrier` (triton::nvidia_gpu::InitBarrierOp)¶

在给定的共享内存分配中初始化一个屏障。

语法：

operation ::= `ttng.init_barrier` $alloc `,` $count attr-dict `:` qualified(type($alloc))

初始化带有内存屏障信息的共享内存分配。 alloc 是共享内存分配的描述符。count 是屏障期望的到达次数。

这会降低到 PTX mbarrier.init.shared::cta.b64。

特性: VerifyTensorLayoutsTrait

属性:¶

属性	MLIR类型	描述
`count`	::mlir::IntegerAttr	32-bit signless integer attribute

操作数:¶

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

`ttng.inval_barrier` (triton::nvidia_gpu::InvalBarrierOp)¶

使屏障分配无效。

语法：

operation ::= `ttng.inval_barrier` $alloc attr-dict `:` qualified(type($alloc))

使屏障分配无效以便可以重新使用。根据PTX规范，必须在重用mbarrier使用的内存之前完成此操作。

https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier-inval

特性: VerifyTensorLayoutsTrait

操作数:¶

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

`ttng.tc_gen5_mma` (triton::nvidia_gpu::TCGen5MMAOp)¶

块级操作映射到tensorcore gen5矩阵乘法累加单元

语法：

operation ::= `ttng.tc_gen5_mma` $a`,` $b`,` $d`,` $useD`,` $pred
              `` custom<BarriersAndPreds>($barriers, $barrier_preds)
              attr-dict `:` qualified(type($a)) `,` qualified(type($b)) `,`
              qualified(type($d)) (`,` qualified(type($barriers))^)?

$d += matrix_multiply($a, $b). 如果没有给定屏障，该操作将被视为同步操作，否则该操作将在给定屏障上触发提交/到达。如果存在屏障，在屏障等待后读取结果是安全的。如果设置了$two_ctas，该操作将在两个连续的CTA之间执行矩阵乘法，它将读取分布在两个CTA之间的数据。如果操作是同步的，还会同步这两个CTA。

特性: SameVariadicOperandSize, VerifyTensorLayoutsTrait

接口: DotOpInterface, MMAv5OpInterface, MemoryEffectOpInterface

属性:¶

属性	MLIR类型	描述
`two_ctas`	::mlir::UnitAttr	unit attribute

操作数:¶

操作数	描述
`a`	Triton IR类型系统中的内存描述符类型 (`::mlir::triton::gpu::MemDescType`)
`b`	Triton IR类型系统中的内存描述符类型 (`::mlir::triton::gpu::MemDescType`)
`d`	Triton IR类型系统中的内存描述符类型 (`::mlir::triton::gpu::MemDescType`)
`useD`	1位无符号整数
`pred`	1位无符号整数
`barriers`	在Triton IR类型系统中，内存描述符类型的可变参数(`::mlir::triton::gpu::MemDescType`)
`barrier_preds`	1位无符号整数的可变参数

`ttng.tc_gen5_mma_scaled` (triton::nvidia_gpu::TCGen5MMAScaledOp)¶

块级操作映射到tensorcore gen5矩阵乘法累加单元

语法：

operation ::= `ttng.tc_gen5_mma_scaled` $a `,` $b `,` $d `,` $a_scale `,` $b_scale `,` $useD`,` $pred
              `lhs` `=` $a_type `rhs` `=` $b_type
              `` custom<BarriersAndPreds>($barriers, $barrier_preds)
              attr-dict `:` qualified(type($a)) `,` qualified(type($b)) `,`
              qualified(type($d)) `,` qualified(type($a_scale)) `,`
              qualified(type($b_scale)) (`,` qualified(type($barriers))^)?

$d += matrix_multiply(scale($lhs, $lhs_scale), scale(rlhs, $rhs_scale)) 如果没有给定屏障，则该操作假定为同步操作；否则该操作将在给定屏障上触发提交/到达。如果存在屏障，则在屏障等待后读取结果是安全的。

特性: SameVariadicOperandSize, VerifyTensorLayoutsTrait

接口: DotOpInterface, MMAv5OpInterface, MemoryEffectOpInterface

属性:¶

属性	MLIR类型	描述
`a_type`	::mlir::triton::ScaleDotElemTypeAttr	allowed 32-bit signless integer cases: 0, 1, 2, 3, 4, 5, 6
`b_type`	::mlir::triton::ScaleDotElemTypeAttr	allowed 32-bit signless integer cases: 0, 1, 2, 3, 4, 5, 6

操作数:¶

操作数	描述
`a`	Triton IR类型系统中的内存描述符类型 (`::mlir::triton::gpu::MemDescType`)
`b`	Triton IR类型系统中的内存描述符类型 (`::mlir::triton::gpu::MemDescType`)
`d`	Triton IR类型系统中的内存描述符类型 (`::mlir::triton::gpu::MemDescType`)
`a_scale`	Triton IR类型系统中的内存描述符类型 (`::mlir::triton::gpu::MemDescType`)
`b_scale`	Triton IR类型系统中的内存描述符类型 (`::mlir::triton::gpu::MemDescType`)
`useD`	1位无符号整数
`pred`	1位无符号整数
`barriers`	内存描述符类型的可变参数 (`::mlir::triton::gpu::MemDescType`) 在Triton IR类型系统中
`barrier_preds`	1位无符号整数的可变参数

`ttng.async_tma_store_wait` (triton::nvidia_gpu::TMAStoreWaitOp)¶

等待所有输入读取完成。

语法：

operation ::= `ttng.async_tma_store_wait` attr-dict

等待所有从关联存储操作中读取的操作完成。在共享内存可以被写入之前，这是必需的。

特性: VerifyTensorLayoutsTrait

属性:¶

属性	MLIR类型	描述
`pendings`	::mlir::IntegerAttr	32-bit signless integer attribute

`ttng.tmem_alloc` (triton::nvidia_gpu::TMEMAllocOp)¶

分配张量内存

语法：

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

此操作在张量内存中分配缓冲区，并返回一个包含地址和缓冲区视图的描述符。这与ttg.local_alloc类似，不同之处在于缓冲区是在张量内存中分配的。

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

特性: VerifyTensorLayoutsTrait

接口: MemoryEffectOpInterface

操作数:¶

操作数	描述
`src`	浮点数、整数或指针值的排序张量

结果:¶

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

`ttng.tmem_copy` (triton::nvidia_gpu::TMEMCopyOp)¶

启动从共享内存到张量内存的异步复制操作。

语法：

operation ::= `ttng.tmem_copy` $src `,` $dst `,` $barrier attr-dict `:` functional-type(operands, results)

连续存储在SMEM中的2D块会按照目标地址指定的方式复制到TMEM中。可以通过等待可选屏障来观察复制操作的完成情况。如果该操作与MMA操作一起使用，一个屏障可以同时用于等待复制和MMA操作完成。我们不需要在MMA操作之前等待复制完成，因为tcgen05.cp后接tcgen05.mma的操作顺序是有保证的。

该操作会转换为PTX指令tcgen05.cp。目前我们仅支持1CTA和该指令的warpx4.32x128b变体。SMEM中的每个32x128b块会在4个warp上进行复制，并存储到TMEM的128行4列中。该操作的主要使用场景是将分块的比例因子从SMEM复制到TMEM。

输入SMEM的形状可以根据使用场景灵活选择。在最简单的情况下（例如单元测试），对于8位值的复制，源SMEM的形状可以是(32 x num_blocks, 16)，而目标TMEM的形状应为(128, 16 x num_blocks)。对于缩放后的GEMM，需要在SMEM中存储rep_m x rep_k个32x128b块的副本，其中rep_m = BLOCK_M / 128，rep_k = BLOCK_K / scale_vec_size / 4，而对于MXFP，scale_vec_size = 32。从概念上讲，SMEM以高维布局(rep_m, rep_k, 32, 4, 4B)组织。可以将某些轴展平为一个，以减少加载的秩。例如，支持以下模式：

(rep_m, rep_k * 32 x 4 x 4B), 使用cp.async进行2D缩放加载
(rep_m, rep_k, 32, 16B), 使用TMA进行4D规模加载
(rep_m, rep_k, 32, 4, 4B)，使用cp.async进行5D缩放加载由于rep_m块在SMEM中不连续，因此该轴无法展平为内部轴。

在Triton中，用于分块缩放（blocked scales）的TMEM内存描述符（memdesc）必须符合以下形式：

它的形状必须为(BLOCK_MN, BLOCK_K / scale_vec_size)，表示分块比例的逻辑形状。
必须与tensor_memory_scales_encoding一起使用，以指示基于块的布局及其在4个warp上的重复。

相比之下，src SMEM必须采用上文描述的显式基于块(chunk-based)的布局。因此，中间表示(IR)可能如下所示：

ttng.tmem_copy %1, %0 : (!ttg.memdesc<1x1x32x4x4xi8, #shared1, #smem>, !ttg.memdesc<128x4xi8, #tmem_scales, #ttng.tensor_memory>) -> ()

我们对此复制操作的语义解释如下。SMEM中基于块的布局意味着TMEM中的逻辑形状(BLOCK_MN, BLOCK_K / scale_vec_size)是某些重塑和转置操作的结果。实际上，为了利用原生缩放布局和TMEM复制操作的优势，用户需要在将scales输入dot_scaled之前执行scales5D.trans(0, 3, 2, 1, 4).reshape(BLOCK_M, BLOCK_K // scale_vec_size)。当我们在IR中使用tmem_copy时，这些重塑和转置操作会被移除。但它们对寄存器造成的逻辑形状变化现在被理解为已融入tmem_copy本身。理想情况下，我们会将对寄存器执行的重塑/转置操作提升到SMEM内存描述符上，使tmem_copy成为一个简单的2D复制操作：(BLOCK_MN, BLOCK_K / scale_vec_size) -> (BLOCK_MN, BLOCK_K / scale_vec_size)。在内存描述符上缺乏此类操作的情况下，我们选择在tmem_copy中隐式编码重塑/转置语义。

特性: VerifyTensorLayoutsTrait

操作数:¶

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

`ttng.tmem_load` (triton::nvidia_gpu::TMEMLoadOp)¶

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

语法：

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

这与ttg.local_load类似，只是结果布局被限制为仅有的几种可能性。因此，我们无法将此操作与任何像local_load这样的布局转换操作结合使用。

特性: VerifyTensorLayoutsTrait

操作数:¶

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

结果:¶

结果	描述
`result`	浮点数、整数或指针值的排序张量

`ttng.tmem_store` (triton::nvidia_gpu::TMEMStoreOp)¶

将一个分布式张量存储到张量内存的缓冲区中

语法：

operation ::= `ttng.tmem_store` $src `,` $dst `,` $pred attr-dict `:` type($src) `->` qualified(type($dst))

这与ttg.local_local类似，只是源布局被限制为仅少数几种可能性。

特性: VerifyTensorLayoutsTrait

操作数:¶

操作数	描述
`dst`	Triton IR类型系统中的内存描述符类型 (`::mlir::triton::gpu::MemDescType`)
`src`	浮点数、整数或指针值的排序张量
`pred`	1位无符号整数

`ttng.tmem_subslice` (triton::nvidia_gpu::TMEMSubSliceOp)¶

获取张量内存分配的子切片

语法：

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

该操作获取张量内存分配的一个子切片，并返回一个新的描述符，包含该子切片的地址和视图。这与ttg.memdesc_subview类似，区别在于偏移量需要是静态的，并且我们只能沿着2D内存描述符的内层维度进行切片，因为这是TMem唯一支持的操作。

特性: AlwaysSpeculatableImplTrait, VerifyTensorLayoutsTrait

接口: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

效果: MemoryEffects::Effect{}

属性:¶

属性	MLIR类型	描述
`N`	::mlir::IntegerAttr	32-bit signless integer attribute

操作数:¶

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

结果:¶

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

`ttng.tensor_desc_to_tma_ptr` (triton::nvidia_gpu::TensorDescToTMAPtrOp)¶

将张量描述符转换为指向tma描述符的指针

语法：

operation ::= `ttng.tensor_desc_to_tma_ptr` $desc attr-dict `:` qualified(type($desc)) `to` qualified(type($ptr))

特性: AlwaysSpeculatableImplTrait, VerifyTensorLayoutsTrait

接口: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

效果: MemoryEffects::Effect{}

操作数:¶

操作数	描述
`描述`	Triton IR类型系统中的张量描述符类型 (`::mlir::triton::TensorDescType`)

结果:¶

结果	描述
`ptr`	ptr

`ttng.wait_barrier` (triton::nvidia_gpu::WaitBarrierOp)¶

等待mbarrier阶段完成。

语法：

operation ::= `ttng.wait_barrier` $alloc `,` $phase (`,` $pred^)? (`deps` $deps^)? attr-dict `:` qualified(type($alloc)) (`,` type($deps)^)?

阻塞程序进度，直到alloc中的mbarrier对象完成其当前阶段。

这通过使用PTX指令mbarrier.try_wait.parity.shared.b64来降低等待循环。

接受可选的内存列表。如果存在，则假定在屏障完成之前可以访问任何依赖项。

The barrier behavior is described here: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-asynchronous-copy-completion-mechanisms

特性: AttrSizedOperandSegments, VerifyTensorLayoutsTrait

操作数:¶

操作数	描述
`alloc`	Triton IR类型系统中的内存描述符类型 (`::mlir::triton::gpu::MemDescType`)
`phase`	32位无符号整数
`pred`	1位无符号整数
`deps`	内存描述符类型的可变参数 (`::mlir::triton::gpu::MemDescType`) 在Triton IR类型系统中

`ttng.warp_group_dot` (triton::nvidia_gpu::WarpGroupDotOp)¶

Warp组点积

语法：

operation ::= `ttng.warp_group_dot` $a`,` $b`,` $c (`,` $useC^)? attr-dict `:` type($a) `*` type($b) `->` type($d)

$d = matrix_multiply($a, $b) + $c。关于InputPrecisionAttr的文档，请参阅TT_DotOp

特性: VerifyTensorLayoutsTrait

接口: DotOpInterface, InferTypeOpInterface, MemoryEffectOpInterface

属性:¶

属性	MLIR类型	描述
`inputPrecision`	::mlir::triton::InputPrecisionAttr	allowed 32-bit signless integer cases: 0, 1, 2
`maxNumImpreciseAcc`	::mlir::IntegerAttr	32-bit signless integer attribute
`isAsync`	::mlir::BoolAttr	bool attribute

操作数:¶

操作数	描述
`a`	TensorOrMemDesc 实例
`b`	TensorOrMemDesc 实例
`c`	浮点或整数值的排名张量
`useC`	1位无符号整数

结果:¶

结果	描述
`d`	浮点或整数值的排序张量

`ttng.warp_group_dot_wait` (triton::nvidia_gpu::WarpGroupDotWaitOp)¶

Warp组点等待

语法：

operation ::= `ttng.warp_group_dot_wait` $inputs attr-dict `:` type($inputs)

等待直到未完成的异步点操作数量不超过$pendings。

$inputs必须是对应于我们正在等待的异步点积操作的张量。例如，如果有N个待处理的异步点积操作，而我们调用warp_group_dot_wait 1，那么$inputs必须是第一个点积操作的结果。

特性: VerifyTensorLayoutsTrait

接口: InferTypeOpInterface

属性:¶

属性	MLIR类型	描述
`pendings`	::mlir::IntegerAttr	32-bit signless integer attribute

操作数:¶

操作数	描述
`inputs`	TensorOrMemDesc 实例的可变参数

结果:¶

结果	描述
`outputs`	TensorOrMemDesc 实例的可变参数

TritonNvidiaGPUOps¶

ttng.arrive_barrier (triton::nvidia_gpu::ArriveBarrierOp)¶

属性:¶

操作数:¶

ttng.async_tma_copy_global_to_local (triton::nvidia_gpu::AsyncTMACopyGlobalToLocalOp)¶

属性:¶

操作数:¶

ttng.async_tma_copy_local_to_global (triton::nvidia_gpu::AsyncTMACopyLocalToGlobalOp)¶

操作数:¶

ttng.async_tma_gather (triton::nvidia_gpu::AsyncTMAGatherOp)¶

操作数:¶

ttng.async_tma_scatter (triton::nvidia_gpu::AsyncTMAScatterOp)¶

操作数:¶

ttng.barrier_expect (triton::nvidia_gpu::BarrierExpectOp)¶

属性:¶

操作数:¶

ttng.cluster_arrive (triton::nvidia_gpu::ClusterArriveOp)¶

属性:¶

ttng.cluster_wait (triton::nvidia_gpu::ClusterWaitOp)¶

ttng.fence_async_shared (triton::nvidia_gpu::FenceAsyncSharedOp)¶

属性:¶

ttng.init_barrier (triton::nvidia_gpu::InitBarrierOp)¶

属性:¶

操作数:¶

ttng.inval_barrier (triton::nvidia_gpu::InvalBarrierOp)¶

操作数:¶

ttng.tc_gen5_mma (triton::nvidia_gpu::TCGen5MMAOp)¶

属性:¶

操作数:¶

ttng.tc_gen5_mma_scaled (triton::nvidia_gpu::TCGen5MMAScaledOp)¶

属性:¶

操作数:¶

ttng.async_tma_store_wait (triton::nvidia_gpu::TMAStoreWaitOp)¶

属性:¶

ttng.tmem_alloc (triton::nvidia_gpu::TMEMAllocOp)¶

操作数:¶

结果:¶

ttng.tmem_copy (triton::nvidia_gpu::TMEMCopyOp)¶

操作数:¶

ttng.tmem_load (triton::nvidia_gpu::TMEMLoadOp)¶

操作数:¶

结果:¶

ttng.tmem_store (triton::nvidia_gpu::TMEMStoreOp)¶

操作数:¶

ttng.tmem_subslice (triton::nvidia_gpu::TMEMSubSliceOp)¶

属性:¶

操作数:¶

结果:¶

ttng.tensor_desc_to_tma_ptr (triton::nvidia_gpu::TensorDescToTMAPtrOp)¶

操作数:¶

结果:¶

ttng.wait_barrier (triton::nvidia_gpu::WaitBarrierOp)¶

操作数:¶

ttng.warp_group_dot (triton::nvidia_gpu::WarpGroupDotOp)¶

属性:¶

操作数:¶

结果:¶

ttng.warp_group_dot_wait (triton::nvidia_gpu::WarpGroupDotWaitOp)¶

属性:¶

操作数:¶

结果:¶

`ttng.arrive_barrier` (triton::nvidia_gpu::ArriveBarrierOp)¶

`ttng.async_tma_copy_global_to_local` (triton::nvidia_gpu::AsyncTMACopyGlobalToLocalOp)¶

`ttng.async_tma_copy_local_to_global` (triton::nvidia_gpu::AsyncTMACopyLocalToGlobalOp)¶

`ttng.async_tma_gather` (triton::nvidia_gpu::AsyncTMAGatherOp)¶

`ttng.async_tma_scatter` (triton::nvidia_gpu::AsyncTMAScatterOp)¶

`ttng.barrier_expect` (triton::nvidia_gpu::BarrierExpectOp)¶

`ttng.cluster_arrive` (triton::nvidia_gpu::ClusterArriveOp)¶

`ttng.cluster_wait` (triton::nvidia_gpu::ClusterWaitOp)¶

`ttng.fence_async_shared` (triton::nvidia_gpu::FenceAsyncSharedOp)¶

`ttng.init_barrier` (triton::nvidia_gpu::InitBarrierOp)¶

`ttng.inval_barrier` (triton::nvidia_gpu::InvalBarrierOp)¶

`ttng.tc_gen5_mma` (triton::nvidia_gpu::TCGen5MMAOp)¶

`ttng.tc_gen5_mma_scaled` (triton::nvidia_gpu::TCGen5MMAScaledOp)¶

`ttng.async_tma_store_wait` (triton::nvidia_gpu::TMAStoreWaitOp)¶

`ttng.tmem_alloc` (triton::nvidia_gpu::TMEMAllocOp)¶

`ttng.tmem_copy` (triton::nvidia_gpu::TMEMCopyOp)¶

`ttng.tmem_load` (triton::nvidia_gpu::TMEMLoadOp)¶

`ttng.tmem_store` (triton::nvidia_gpu::TMEMStoreOp)¶

`ttng.tmem_subslice` (triton::nvidia_gpu::TMEMSubSliceOp)¶

`ttng.tensor_desc_to_tma_ptr` (triton::nvidia_gpu::TensorDescToTMAPtrOp)¶

`ttng.wait_barrier` (triton::nvidia_gpu::WaitBarrierOp)¶

`ttng.warp_group_dot` (triton::nvidia_gpu::WarpGroupDotOp)¶

`ttng.warp_group_dot_wait` (triton::nvidia_gpu::WarpGroupDotWaitOp)¶