NVGPUOps¶

`nvgpu.cga_barrier_arrive` (::mlir::triton::nvgpu::CGABarrierArriveOp)¶

Syntax:

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

`nvgpu.cga_barrier_sync` (::mlir::triton::nvgpu::CGABarrierSyncOp)¶

Syntax:

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

`nvgpu.cga_barrier_wait` (::mlir::triton::nvgpu::CGABarrierWaitOp)¶

Syntax:

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

`nvgpu.cluster_arrive` (::mlir::triton::nvgpu::ClusterArriveOp)¶

Syntax:

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

Attributes:¶

Attribute	MLIR Type	Description
`relaxed`	::mlir::IntegerAttr	1-bit signless integer attribute

`nvgpu.cluster_id` (::mlir::triton::nvgpu::ClusterCTAIdOp)¶

Syntax:

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

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Results:¶

Result	Description
`result`	32-bit signless integer

`nvgpu.cluster_wait` (::mlir::triton::nvgpu::ClusterWaitOp)¶

Syntax:

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

`nvgpu.fence_async_shared` (::mlir::triton::nvgpu::FenceAsyncSharedOp)¶

Syntax:

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

Attributes:¶

Attribute	MLIR Type	Description
`bCluster`	::mlir::BoolAttr	bool attribute

`nvgpu.fence_mbarrier_init` (::mlir::triton::nvgpu::FenceMBarrierInitOp)¶

Syntax:

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

`nvgpu.load_dsmem` (::mlir::triton::nvgpu::LoadDSmemOp)¶

Syntax:

operation ::= `nvgpu.load_dsmem` operands attr-dict `:` functional-type(operands, results)

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

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

Attributes:¶

Attribute	MLIR Type	Description
`bitwidth`	::mlir::IntegerAttr	32-bit signless integer attribute
`vec`	::mlir::IntegerAttr	32-bit signless integer attribute

Operands:¶

Operand	Description
`addr`	LLVM pointer type
`ctaId`	32-bit signless integer

Results:¶

Result	Description
`result`	LLVM type with size

`nvgpu.mbarrier_arrive` (::mlir::triton::nvgpu::MBarrierArriveOp)¶

Syntax:

operation ::= `nvgpu.mbarrier_arrive` $mbarrier `,` $pred (`,` $ctaId^)? attr-dict `:` type($mbarrier)

Attributes:¶

Attribute	MLIR Type	Description
`arriveType`	::mlir::triton::nvgpu::MBarriveTypeAttr	mbarrier arrive type, either ‘normal’, ‘expect_tx’, ‘cp_async’
`txCount`	::mlir::IntegerAttr	32-bit signless integer attribute

Operands:¶

Operand	Description
`mbarrier`	LLVM pointer to 64-bit signless integer
`pred`	1-bit signless integer
`ctaId`	32-bit signless integer

`nvgpu.mbarrier_init` (::mlir::triton::nvgpu::MBarrierInitOp)¶

Syntax:

operation ::= `nvgpu.mbarrier_init` $mbarrier `,` $pred attr-dict `:` type($mbarrier)

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

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

Attributes:¶

Attribute	MLIR Type	Description
`count`	::mlir::IntegerAttr	32-bit signless integer attribute

Operands:¶

Operand	Description
`mbarrier`	LLVM pointer to 64-bit signless integer
`pred`	1-bit signless integer

`nvgpu.mbarrier_wait` (::mlir::triton::nvgpu::MBarrierWaitOp)¶

Syntax:

operation ::= `nvgpu.mbarrier_wait` $mbarrier `,` $phase attr-dict `:` type(operands)

Operands:¶

Operand	Description
`mbarrier`	LLVM pointer to 64-bit signless integer
`phase`	1-bit signless integer

`nvgpu.bar_arrive` (::mlir::triton::nvgpu::NamedBarrierArriveOp)¶

Syntax:

operation ::= `nvgpu.bar_arrive` $bar `,` $numThreads attr-dict `:` type(operands)

Operands:¶

Operand	Description
`bar`	32-bit signless integer
`numThreads`	32-bit signless integer

`nvgpu.bar_wait` (::mlir::triton::nvgpu::NamedBarrierWaitOp)¶

Syntax:

operation ::= `nvgpu.bar_wait` $bar `,` $numThreads attr-dict `:` type(operands)

Operands:¶

Operand	Description
`bar`	32-bit signless integer
`numThreads`	32-bit signless integer

`nvgpu.offset_of_stmatrix_v4` (::mlir::triton::nvgpu::OffsetOfStmatrixV4Op)¶

Syntax:

operation ::= `nvgpu.offset_of_stmatrix_v4` operands attr-dict `:` type(operands) `->` type($offset)

Attributes:¶

Attribute	MLIR Type	Description
`leadingDimOffset`	::mlir::IntegerAttr	32-bit signless integer attribute
`rowStride`	::mlir::IntegerAttr	32-bit signless integer attribute
`swizzleEnabled`	::mlir::IntegerAttr	1-bit signless integer attribute

Operands:¶

Operand	Description
`threadId`	32-bit signless integer
`rowOfWarp`	32-bit signless integer
`elemIdx`	32-bit signless integer

Results:¶

Result	Description
`offset`	32-bit signless integer

`nvgpu.offset_of_sts64` (::mlir::triton::nvgpu::OffsetOfSts64Op)¶

Syntax:

operation ::= `nvgpu.offset_of_sts64` operands attr-dict `:` type(operands) `->` type($offset)

Attributes:¶

Attribute	MLIR Type	Description
`leadingDimOffset`	::mlir::IntegerAttr	32-bit signless integer attribute
`rowStride`	::mlir::IntegerAttr	32-bit signless integer attribute
`swizzleEnabled`	::mlir::IntegerAttr	1-bit signless integer attribute

Operands:¶

Operand	Description
`threadId`	32-bit signless integer
`rowOfWarp`	32-bit signless integer
`elemIdx`	32-bit signless integer

Results:¶

Result	Description
`offset`	32-bit signless integer

`nvgpu.reg_alloc` (::mlir::triton::nvgpu::RegAllocOp)¶

Syntax:

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

Attributes:¶

Attribute	MLIR Type	Description
`regCount`	::mlir::IntegerAttr	32-bit signless integer attribute

`nvgpu.reg_dealloc` (::mlir::triton::nvgpu::RegDeallocOp)¶

Syntax:

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

Attributes:¶

Attribute	MLIR Type	Description
`regCount`	::mlir::IntegerAttr	32-bit signless integer attribute

`nvgpu.store_dsmem` (::mlir::triton::nvgpu::StoreDSmemOp)¶

Syntax:

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

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

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

Operands:¶

Operand	Description
`addr`	LLVM pointer type
`ctaId`	32-bit signless integer
`values`	LLVM type with size
`pred`	1-bit signless integer

`nvgpu.stmatrix` (::mlir::triton::nvgpu::StoreMatrixOp)¶

Syntax:

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

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

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

Operands:¶

Operand	Description
`addr`	LLVM pointer to 8-bit signless integer
`datas`	32-bit signless integer

`nvgpu.sts64` (::mlir::triton::nvgpu::Sts64Op)¶

Syntax:

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

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

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

Operands:¶

Operand	Description
`offset`	32-bit signless integer
`d0`	32-bit float or 32-bit signless integer
`d1`	32-bit float or 32-bit signless integer

`nvgpu.tma_load_im2col` (::mlir::triton::nvgpu::TMALoadIm2colOp)¶

Syntax:

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

Attributes:¶

Attribute	MLIR Type	Description
`mcastMask`	::mlir::IntegerAttr	16-bit signless integer attribute

Operands:¶

Operand	Description
`dst`	LLVM pointer to 8-bit signless integer
`mbarrier`	LLVM pointer to 64-bit signless integer
`tmaDesc`	LLVM pointer to 8-bit signless integer
`l2Desc`	64-bit signless integer
`im2colOffsets`	LLVM structure type
`pred`	1-bit signless integer
`coords`	32-bit signless integer

`nvgpu.tma_load_tiled` (::mlir::triton::nvgpu::TMALoadTiledOp)¶

Syntax:

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

Traits: AttrSizedOperandSegments

Operands:¶

Operand	Description
`dst`	LLVM pointer to 8-bit signless integer
`mbarrier`	LLVM pointer to 64-bit signless integer
`tmaDesc`	LLVM pointer to 8-bit signless integer
`l2Desc`	64-bit signless integer
`pred`	1-bit signless integer
`coords`	32-bit signless integer
`mcastMask`	16-bit signless integer

`nvgpu.tma_store_tiled` (::mlir::triton::nvgpu::TMAStoreTiledOp)¶

Syntax:

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

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

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

Operands:¶

Operand	Description
`tmaDesc`	LLVM pointer to 8-bit signless integer
`src`	LLVM pointer to 8-bit signless integer
`pred`	1-bit signless integer
`coords`	32-bit signless integer

`nvgpu.wgmma_commit_group` (::mlir::triton::nvgpu::WGMMACommitGroupOp)¶

Syntax:

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

`nvgpu.wgmma_desc_create` (::mlir::triton::nvgpu::WGMMADescCreateOp)¶

Syntax:

operation ::= `nvgpu.wgmma_desc_create` $buffer `,` $height attr-dict `:` functional-type(operands, results)

Attributes:¶

Attribute	MLIR Type	Description
`mode`	::mlir::triton::nvgpu::WGMMADescModeAttr	wgmma desc mode, either ‘none’, ‘swizzle128’, ‘swizzle64’, or ‘swizzle32’
`swizzling`	::mlir::IntegerAttr	64-bit signless integer attribute

Operands:¶

Operand	Description
`buffer`	LLVM pointer type
`height`	32-bit signless integer

Results:¶

Result	Description
`res`	64-bit signless integer

`nvgpu.wgmma_fence` (::mlir::triton::nvgpu::WGMMAFenceOp)¶

Syntax:

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

`nvgpu.wgmma` (::mlir::triton::nvgpu::WGMMAOp)¶

Syntax:

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

Attributes:¶

Attribute	MLIR Type	Description
`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’

Operands:¶

Operand	Description
`opA`	wgmma operand A/B type
`opB`	wgmma operand A/B type
`opC`	LLVM structure type

Results:¶

Result	Description
`res`	LLVM structure type

`nvgpu.wgmma_wait_group` (::mlir::triton::nvgpu::WGMMAWaitGroupOp)¶

Syntax:

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

Attributes:¶

Attribute	MLIR Type	Description
`pendings`	::mlir::IntegerAttr	32-bit signless integer attribute

NVGPUOps¶

nvgpu.cga_barrier_arrive (::mlir::triton::nvgpu::CGABarrierArriveOp)¶

nvgpu.cga_barrier_sync (::mlir::triton::nvgpu::CGABarrierSyncOp)¶

nvgpu.cga_barrier_wait (::mlir::triton::nvgpu::CGABarrierWaitOp)¶

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

Attributes:¶

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

Results:¶

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

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

Attributes:¶

nvgpu.fence_mbarrier_init (::mlir::triton::nvgpu::FenceMBarrierInitOp)¶

nvgpu.load_dsmem (::mlir::triton::nvgpu::LoadDSmemOp)¶

Attributes:¶

Operands:¶

Results:¶

nvgpu.mbarrier_arrive (::mlir::triton::nvgpu::MBarrierArriveOp)¶

Attributes:¶

Operands:¶

nvgpu.mbarrier_init (::mlir::triton::nvgpu::MBarrierInitOp)¶

Attributes:¶

Operands:¶

nvgpu.mbarrier_wait (::mlir::triton::nvgpu::MBarrierWaitOp)¶

Operands:¶

nvgpu.bar_arrive (::mlir::triton::nvgpu::NamedBarrierArriveOp)¶

Operands:¶

nvgpu.bar_wait (::mlir::triton::nvgpu::NamedBarrierWaitOp)¶

Operands:¶

nvgpu.offset_of_stmatrix_v4 (::mlir::triton::nvgpu::OffsetOfStmatrixV4Op)¶

Attributes:¶

Operands:¶

Results:¶

nvgpu.offset_of_sts64 (::mlir::triton::nvgpu::OffsetOfSts64Op)¶

Attributes:¶

Operands:¶

Results:¶

nvgpu.reg_alloc (::mlir::triton::nvgpu::RegAllocOp)¶

Attributes:¶

nvgpu.reg_dealloc (::mlir::triton::nvgpu::RegDeallocOp)¶

Attributes:¶

nvgpu.store_dsmem (::mlir::triton::nvgpu::StoreDSmemOp)¶

Operands:¶

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

Operands:¶

nvgpu.sts64 (::mlir::triton::nvgpu::Sts64Op)¶

Operands:¶

nvgpu.tma_load_im2col (::mlir::triton::nvgpu::TMALoadIm2colOp)¶

Attributes:¶

Operands:¶

nvgpu.tma_load_tiled (::mlir::triton::nvgpu::TMALoadTiledOp)¶

Operands:¶

nvgpu.tma_store_tiled (::mlir::triton::nvgpu::TMAStoreTiledOp)¶

Operands:¶

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

nvgpu.wgmma_desc_create (::mlir::triton::nvgpu::WGMMADescCreateOp)¶

Attributes:¶

Operands:¶

Results:¶

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

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

Attributes:¶

Operands:¶

Results:¶

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

Attributes:¶

`nvgpu.cga_barrier_arrive` (::mlir::triton::nvgpu::CGABarrierArriveOp)¶

`nvgpu.cga_barrier_sync` (::mlir::triton::nvgpu::CGABarrierSyncOp)¶

`nvgpu.cga_barrier_wait` (::mlir::triton::nvgpu::CGABarrierWaitOp)¶

`nvgpu.cluster_arrive` (::mlir::triton::nvgpu::ClusterArriveOp)¶

`nvgpu.cluster_id` (::mlir::triton::nvgpu::ClusterCTAIdOp)¶

`nvgpu.cluster_wait` (::mlir::triton::nvgpu::ClusterWaitOp)¶

`nvgpu.fence_async_shared` (::mlir::triton::nvgpu::FenceAsyncSharedOp)¶

`nvgpu.fence_mbarrier_init` (::mlir::triton::nvgpu::FenceMBarrierInitOp)¶

`nvgpu.load_dsmem` (::mlir::triton::nvgpu::LoadDSmemOp)¶

`nvgpu.mbarrier_arrive` (::mlir::triton::nvgpu::MBarrierArriveOp)¶

`nvgpu.mbarrier_init` (::mlir::triton::nvgpu::MBarrierInitOp)¶

`nvgpu.mbarrier_wait` (::mlir::triton::nvgpu::MBarrierWaitOp)¶

`nvgpu.bar_arrive` (::mlir::triton::nvgpu::NamedBarrierArriveOp)¶

`nvgpu.bar_wait` (::mlir::triton::nvgpu::NamedBarrierWaitOp)¶

`nvgpu.offset_of_stmatrix_v4` (::mlir::triton::nvgpu::OffsetOfStmatrixV4Op)¶

`nvgpu.offset_of_sts64` (::mlir::triton::nvgpu::OffsetOfSts64Op)¶

`nvgpu.reg_alloc` (::mlir::triton::nvgpu::RegAllocOp)¶

`nvgpu.reg_dealloc` (::mlir::triton::nvgpu::RegDeallocOp)¶

`nvgpu.store_dsmem` (::mlir::triton::nvgpu::StoreDSmemOp)¶

`nvgpu.stmatrix` (::mlir::triton::nvgpu::StoreMatrixOp)¶

`nvgpu.sts64` (::mlir::triton::nvgpu::Sts64Op)¶

`nvgpu.tma_load_im2col` (::mlir::triton::nvgpu::TMALoadIm2colOp)¶

`nvgpu.tma_load_tiled` (::mlir::triton::nvgpu::TMALoadTiledOp)¶

`nvgpu.tma_store_tiled` (::mlir::triton::nvgpu::TMAStoreTiledOp)¶

`nvgpu.wgmma_commit_group` (::mlir::triton::nvgpu::WGMMACommitGroupOp)¶

`nvgpu.wgmma_desc_create` (::mlir::triton::nvgpu::WGMMADescCreateOp)¶

`nvgpu.wgmma_fence` (::mlir::triton::nvgpu::WGMMAFenceOp)¶

`nvgpu.wgmma` (::mlir::triton::nvgpu::WGMMAOp)¶

`nvgpu.wgmma_wait_group` (::mlir::triton::nvgpu::WGMMAWaitGroupOp)¶