NVWSOps¶
nvws.aref.create (triton::nvws::ArefCreateOp)¶
创建一个异步引用。
语法
operation ::= `nvws.aref.create` $operands attr-dict `:` type($result)
创建一个异步引用。
接受可变数量的操作数作为输入，并返回一个 ARef。输入预期为类数组（即 Tensor、MemDesc 等），且所有输入的形状的第一个轴应匹配，表示值的多缓冲。
特性: AlwaysSpeculatableImplTrait
接口: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)
效果: MemoryEffects::Effect{}
操作数:¶
操作数 |
描述 |
|---|---|
|
任意类型的可变参数 |
结果:¶
结果 |
描述 |
|---|---|
|
异步引用 |
nvws.aref.destroy (triton::nvws::ArefDestroyOp)¶
销毁一个异步引用。
语法
operation ::= `nvws.aref.destroy` $aref attr-dict `:` type($aref)
销毁一个异步引用。
操作数:¶
操作数 |
描述 |
|---|---|
|
异步引用 |
nvws.aref.get.enter (triton::nvws::ArefGetEnterOp)¶
进入 ArefGet 区域，可在其中读取缓冲区数据
语法
operation ::= `nvws.aref.get.enter` $aref `[` $stage `,` $phase `]` attr-dict
`:` type($aref) `->` type($results)
进入一个可以自由读取缓冲区数据的“区域”）。这些 ArefGet“区域”可以跨越多次迭代。
操作数:¶
操作数 |
描述 |
|---|---|
|
异步引用 |
|
32位无符号整数 |
|
32位无符号整数 |
结果:¶
结果 |
描述 |
|---|---|
|
任意类型的可变参数 |
nvws.aref.get.exit (triton::nvws::ArefGetExitOp)¶
退出 ArefGet 区域，此后不应再使用缓冲区
语法
operation ::= `nvws.aref.get.exit` $aref `[` $stage `]` $async_ops attr-dict
`:` type($aref)
离开可以自由读取缓冲区数据的区域）。这些 ArefGet“区域”可以跨越多次迭代。
属性：¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
async_ops | ::mlir::ArrayAttr | 异步操作属性数组 |
操作数:¶
操作数 |
描述 |
|---|---|
|
异步引用 |
|
32位无符号整数 |
nvws.aref.put.enter (triton::nvws::ArefPutEnterOp)¶
进入 ArefPut 区域，可在其中写入缓冲区数据
语法
operation ::= `nvws.aref.put.enter` $aref `[` $stage `,` $phase `]` attr-dict
`:` type($aref) `->` type($results)
进入一个可以自由写入缓冲区数据的“区域”）。这些 ArefPut“区域”可以跨越多次迭代。
操作数:¶
操作数 |
描述 |
|---|---|
|
异步引用 |
|
32位无符号整数 |
|
32位无符号整数 |
结果:¶
结果 |
描述 |
|---|---|
|
任意类型的可变参数 |
nvws.aref.put.exit (triton::nvws::ArefPutExitOp)¶
退出 ArefPut 区域，此后不应再使用缓冲区
语法
operation ::= `nvws.aref.put.exit` $aref `[` $stage `]` $async_ops attr-dict
`:` type($aref)
离开可以自由写入缓冲区数据的区域）。这些 ArefPut“区域”可以跨越多次迭代。
属性:¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
async_ops | ::mlir::ArrayAttr | 异步操作属性数组 |
操作数:¶
操作数 |
描述 |
|---|---|
|
异步引用 |
|
32位无符号整数 |
nvws.consumer_release (triton::nvws::ConsumerReleaseOp)¶
消费者释放令牌
语法
operation ::= `nvws.consumer_release` $token `,` $idx attr-dict `:` type(operands)
消费者将释放令牌，并通知生产者缓冲区已准备好填充。
操作数:¶
操作数 |
描述 |
|---|---|
|
值张量 |
|
32位无符号整数 |
nvws.consumer_wait (triton::nvws::ConsumerWaitOp)¶
消费者等待缓冲区就绪
语法
operation ::= `nvws.consumer_wait` $token `,` $idx `,` $phase attr-dict `:` type(operands)
消费者将等待缓冲区准备好被消费。如果缓冲区尚未准备好，消费者将等待生产者完成缓冲区填充并释放令牌的信号。
操作数:¶
操作数 |
描述 |
|---|---|
|
值张量 |
|
32位无符号整数 |
|
1位无符号整数 |
nvws.create_token (triton::nvws::CreateTokenOp)¶
创建一个用于通信通道同步的令牌
语法
operation ::= `nvws.create_token` attr-dict `:` type($result)
生产者和消费者将使用令牌进行同步。生产者将获取并持有令牌，直到填满缓冲区，并通知等待的消费者。消费者将持有令牌，直到消费完缓冲区，并通知尝试获取令牌的等待中的生产者。
属性:¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
numBuffers | ::mlir::IntegerAttr | 32位无符号整数属性 |
loadType | ::mlir::triton::nvws::TokenLoadTypeAttr | 允许的32位无符号整数情况: 0, 1, 2, 3, 4 |
结果:¶
结果 |
描述 |
|---|---|
|
值张量 |
nvws.producer_acquire (triton::nvws::ProducerAcquireOp)¶
生产者获取令牌以填充缓冲区
语法
operation ::= `nvws.producer_acquire` $token `,` $idx `,` $phase attr-dict `:` type(operands)
生产者将在填充缓冲区之前尝试获取令牌。如果缓冲区尚未准备好填充，生产者将等待消费者完成缓冲区消费并释放令牌的信号。
操作数:¶
操作数 |
描述 |
|---|---|
|
值张量 |
|
32位无符号整数 |
|
1位无符号整数 |
nvws.producer_commit (triton::nvws::ProducerCommitOp)¶
生产者提交缓冲区更改
语法
operation ::= `nvws.producer_commit` $token `,` $idx attr-dict `:` type(operands)
生产者将释放令牌，并通知消费者缓冲区已准备好被消费。
操作数:¶
操作数 |
描述 |
|---|---|
|
值张量 |
|
32位无符号整数 |
nvws.warp_group (triton::nvws::WarpGroupOp)¶
用于 Warp 专门化的容器操作
用于 Warp 专门化分析的高级容器。
包含可变数量的 Warp 组，每组中 Warp 的数量，以及一个用于保存该 Warp 组计算的区域。
此操作的结果（如果有）是第一个区域的结果，由 nvws.warp_group.yield 操作返回。
nvws.warp_group 在执行前应降低为 ttg.warp_specialize。
特性: RecursiveMemoryEffects, RecursivelySpeculatableImplTrait
接口: ConditionallySpeculatable
属性:¶
| 属性 | MLIR 类型 | 描述 |
|---|---|---|
numWarps | ::mlir::DenseI32ArrayAttr | i32 密集数组属性 |
结果:¶
结果 |
描述 |
|---|---|
|
任意类型的可变参数 |
nvws.warp_group.return (triton::nvws::WarpGroupReturnOp)¶
Warp 组区域的终结符
语法
operation ::= `nvws.warp_group.return` attr-dict
Warp 组预期通过对其输入的引用修改来返回值。因此，warp_group.return 操作不接受从 Warp 组返回的任何值。
特性: AlwaysSpeculatableImplTrait, HasParent<WarpGroupOp>, Terminator
接口: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)
效果: MemoryEffects::Effect{}
nvws.warp_group.yield (triton::nvws::WarpGroupYieldOp)¶
从 nvws.warp_group 的第一个区域产出
语法
operation ::= `nvws.warp_group.yield` ($values^)? attr-dict (`:` type($values)^)?
此操作等同于 ttg.warp_specialize 操作的 ttg.warp_yield 操作。
待办：决定是否应将 nvws.warp_group 移动到 TritonGPU，或者继续让 TritonGPU 依赖于 NVWS。如果是前者，此操作可以移除。后者则涉及 TritonGPU 和 NVWS 之间的循环依赖。
特性: AlwaysSpeculatableImplTrait, HasParent<WarpGroupOp>, ReturnLike, Terminator
接口: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface), RegionBranchTerminatorOpInterface
效果: MemoryEffects::Effect{}
操作数:¶
操作数 |
描述 |
|---|---|
|
任意类型的可变参数 |