NVGPUOps

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

语法

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

属性：

属性	MLIR 类型	描述
relaxed	::mlir::IntegerAttr	1位无符号整型属性

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

语法

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

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

Effects: 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	布尔属性

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

语法

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

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: 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.ldmatrix (triton::nvgpu::LoadMatrixOp)

语法

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

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Read on ::mlir::SideEffects::DefaultResource}

属性：

属性	MLIR 类型	描述
trans	::mlir::UnitAttr	单位属性

操作数：

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

结果：

结果	描述
result	LLVM 结构体类型

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

语法

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

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::SideEffects::DefaultResource}

属性：

属性	MLIR 类型	描述
trans	::mlir::UnitAttr	单位属性

操作数：

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

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

语法

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

表示内核中张量内存基地址的 Op。这用于简化从 TritonGPU 到 LLVM 的转换。

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

Effects: 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位无符号整型属性
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

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

结果：

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