NVGPUOps

nvgpu.cluster_arrive (triton::nvgpu::ClusterArriveOp)

语法：

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

属性:

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

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

语法：

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

特性: AlwaysSpeculatableImplTrait

接口: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

效果: MemoryEffects::Effect{}

结果:

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

nvgpu.cluster_wait (triton::nvgpu::ClusterWaitOp)

语法：

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

nvgpu.fence_async_shared (triton::nvgpu::FenceAsyncSharedOp)

语法：

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

属性:

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

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	allowed 32-bit signless integer cases: 1, 2, 3, 4
scope	::mlir::triton::nvgpu::MemSyncScopeAttr	allowed 32-bit signless integer cases: 1, 2, 3

操作数:

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

结果:

结果	描述
result	浮点数或整数

nvgpu.ldmatrix (triton::nvgpu::LoadMatrixOp)

语法：

operation ::= `nvgpu.ldmatrix` $addr attr-dict `:` functional-type($addr, $result)

接口：MemoryEffectOpInterface (MemoryEffectOpInterface)

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

属性:

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

操作数:

操作数	描述
addr	地址空间3中的LLVM指针

结果:

结果	描述
result	LLVM结构类型

nvgpu.stmatrix (triton::nvgpu::StoreMatrixOp)

语法：

operation ::= `nvgpu.stmatrix` operands attr-dict `:` type(operands)

接口: MemoryEffectOpInterface (MemoryEffectOpInterface)

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

属性:

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

操作数:

操作数	描述
addr	地址空间3中的LLVM指针
vals	32位无符号整数的可变参数

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

语法：

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

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

特性: AlwaysSpeculatableImplTrait

接口: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

效果: MemoryEffects::Effect{}

结果:

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

nvgpu.wgmma_commit_group (triton::nvgpu::WGMMACommitGroupOp)

语法：

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

nvgpu.wgmma_fence (triton::nvgpu::WGMMAFenceOp)

语法：

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

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

语法：

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

属性:

属性	MLIR类型	描述
m	::mlir::IntegerAttr	32-bit signless integer attribute
n	::mlir::IntegerAttr	32-bit signless integer attribute
k	::mlir::IntegerAttr	32-bit signless integer attribute
eltTypeC	::mlir::triton::nvgpu::WGMMAEltTypeAttr	wgmma operand type, either 's8', 's32', 'e4m3', 'e5m2', 'f16', 'bf16', 'tf32', or 'f32'
eltTypeA	::mlir::triton::nvgpu::WGMMAEltTypeAttr	wgmma operand type, either 's8', 's32', 'e4m3', 'e5m2', 'f16', 'bf16', 'tf32', or 'f32'
eltTypeB	::mlir::triton::nvgpu::WGMMAEltTypeAttr	wgmma operand type, either 's8', 's32', 'e4m3', 'e5m2', 'f16', 'bf16', 'tf32', or 'f32'
layoutA	::mlir::triton::nvgpu::WGMMALayoutAttr	wgmma layout, either 'row' or 'col'
layoutB	::mlir::triton::nvgpu::WGMMALayoutAttr	wgmma layout, either 'row' or '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)

接口: InferTypeOpInterface

属性:

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

操作数:

操作数	描述
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位无符号整数