NVGPUOps

nvgpu.cluster_id (triton::nvgpu::ClusterCTAIdOp)

语法

operation ::= `nvgpu.cluster_id` attr-dict

特性: AlwaysSpeculatableImplTrait

接口：ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

效果: MemoryEffects::Effect{}

结果:

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

nvgpu.ld_acquire (triton::nvgpu::LoadAcquireOp)

语法

operation ::= `nvgpu.ld_acquire` $sem `,` $scope `,` $addr (`,` $mask^)? attr-dict `:` functional-type($addr, $result)

接口：MemoryEffectOpInterface (MemoryEffectOpInterface)

效果：MemoryEffects::Effect{MemoryEffects::Read on ::mlir::SideEffects::DefaultResource}

属性：

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

操作数:

操作数	描述
addr	地址空间 1 中的 LLVM 指针
mask	1位无符号整数

结果：

结果	描述
result	浮点或整数

nvgpu.tensor_memory_base (triton::nvgpu::TensorMemoryBaseAddress)

语法

operation ::= `nvgpu.tensor_memory_base` attr-dict

用于表示内核中张量内存基地址的操作。这用于简化从 TritonGPU 到 LLVM 的降级。

特性: AlwaysSpeculatableImplTrait

接口：ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

效果: MemoryEffects::Effect{}

结果：

结果	描述
result	地址空间 6 中的 LLVM 指针

nvgpu.wgmma (triton::nvgpu::WGMMAOp)

语法

operation ::= `nvgpu.wgmma` $opA `,` $opB `,` $useC (`,` $opC^)? attr-dict `:` functional-type(operands, $res)

属性 (Attributes):

属性	MLIR 类型	描述
m	::mlir::IntegerAttr	32位无符号整数属性
n	::mlir::IntegerAttr	32位无符号整数属性
k	::mlir::IntegerAttr	32位无符号整数属性
eltTypeC	::mlir::triton::nvgpu::WGMMAEltTypeAttr	wgmma 操作数类型，可以是 's8'、's32'、'e4m3'、'e5m2'、'f16'、'bf16'、'tf32' 或 'f32'
eltTypeA	::mlir::triton::nvgpu::WGMMAEltTypeAttr	wgmma 操作数类型，可以是 's8'、's32'、'e4m3'、'e5m2'、'f16'、'bf16'、'tf32' 或 'f32'
eltTypeB	::mlir::triton::nvgpu::WGMMAEltTypeAttr	wgmma 操作数类型，可以是 's8'、's32'、'e4m3'、'e5m2'、'f16'、'bf16'、'tf32' 或 'f32'
layoutA	::mlir::triton::nvgpu::WGMMALayoutAttr	wgmma 布局，可以是 'row' 或 'col'
layoutB	::mlir::triton::nvgpu::WGMMALayoutAttr	wgmma 布局，可以是 'row' 或 'col'

操作数:

操作数	描述
opA	wgmma 操作数 A/B 类型
opB	wgmma 操作数 A/B 类型
useC	1位无符号整数
opC	LLVM 结构体类型

结果：

结果	描述
res	LLVM 结构体类型

nvgpu.wgmma_wait_group (triton::nvgpu::WGMMAWaitGroupOp)

语法

operation ::= `nvgpu.wgmma_wait_group` $input attr-dict `:` type($input)

接口 (Interfaces): InferTypeOpInterface

属性：

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

操作数：

操作数	描述
input	LLVM 结构体类型

结果：

结果	描述
output	LLVM 结构体类型

nvgpu.warp_id (triton::nvgpu::WarpIdOp)

语法

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

特性: AlwaysSpeculatableImplTrait

接口：ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

效果: MemoryEffects::Effect{}

结果：

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