MLIR

Multi-Level IR Compiler Framework

'nvvm' Dialect

Operation definition 

nvvm.barrier0 (::mlir::NVVM::Barrier0Op) 

Syntax:

operation ::= `nvvm.barrier0` attr-dict

nvvm.read.ptx.sreg.ntid.x (::mlir::NVVM::BlockDimXOp) 

Syntax:

operation ::= `nvvm.read.ptx.sreg.ntid.x` attr-dict `:` type($res)

Interfaces: NoSideEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Results: 

ResultDescription
resLLVM dialect-compatible type

nvvm.read.ptx.sreg.ntid.y (::mlir::NVVM::BlockDimYOp) 

Syntax:

operation ::= `nvvm.read.ptx.sreg.ntid.y` attr-dict `:` type($res)

Interfaces: NoSideEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Results: 

ResultDescription
resLLVM dialect-compatible type

nvvm.read.ptx.sreg.ntid.z (::mlir::NVVM::BlockDimZOp) 

Syntax:

operation ::= `nvvm.read.ptx.sreg.ntid.z` attr-dict `:` type($res)

Interfaces: NoSideEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Results: 

ResultDescription
resLLVM dialect-compatible type

nvvm.read.ptx.sreg.ctaid.x (::mlir::NVVM::BlockIdXOp) 

Syntax:

operation ::= `nvvm.read.ptx.sreg.ctaid.x` attr-dict `:` type($res)

Interfaces: NoSideEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Results: 

ResultDescription
resLLVM dialect-compatible type

nvvm.read.ptx.sreg.ctaid.y (::mlir::NVVM::BlockIdYOp) 

Syntax:

operation ::= `nvvm.read.ptx.sreg.ctaid.y` attr-dict `:` type($res)

Interfaces: NoSideEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Results: 

ResultDescription
resLLVM dialect-compatible type

nvvm.read.ptx.sreg.ctaid.z (::mlir::NVVM::BlockIdZOp) 

Syntax:

operation ::= `nvvm.read.ptx.sreg.ctaid.z` attr-dict `:` type($res)

Interfaces: NoSideEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Results: 

ResultDescription
resLLVM dialect-compatible type

nvvm.read.ptx.sreg.nctaid.x (::mlir::NVVM::GridDimXOp) 

Syntax:

operation ::= `nvvm.read.ptx.sreg.nctaid.x` attr-dict `:` type($res)

Interfaces: NoSideEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Results: 

ResultDescription
resLLVM dialect-compatible type

nvvm.read.ptx.sreg.nctaid.y (::mlir::NVVM::GridDimYOp) 

Syntax:

operation ::= `nvvm.read.ptx.sreg.nctaid.y` attr-dict `:` type($res)

Interfaces: NoSideEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Results: 

ResultDescription
resLLVM dialect-compatible type

nvvm.read.ptx.sreg.nctaid.z (::mlir::NVVM::GridDimZOp) 

Syntax:

operation ::= `nvvm.read.ptx.sreg.nctaid.z` attr-dict `:` type($res)

Interfaces: NoSideEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Results: 

ResultDescription
resLLVM dialect-compatible type

nvvm.read.ptx.sreg.laneid (::mlir::NVVM::LaneIdOp) 

Syntax:

operation ::= `nvvm.read.ptx.sreg.laneid` attr-dict `:` type($res)

Interfaces: NoSideEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Results: 

ResultDescription
resLLVM dialect-compatible type

nvvm.mma.sync (::mlir::NVVM::MmaOp) 

Syntax:

operation ::= `nvvm.mma.sync` $args attr-dict `:` functional-type($args, $res)

Operands: 

OperandDescription
argsLLVM dialect-compatible type

Results: 

ResultDescription
resLLVM dialect-compatible type

nvvm.shfl.sync (::mlir::NVVM::ShflOp) 

Syntax:

operation ::= `nvvm.shfl.sync` $kind $dst `,` $val `,` $offset `,` $mask_and_clamp  attr-dict
              `:` type($val) `->` type($res)

Attributes: 

AttributeMLIR TypeDescription
kindmlir::StringAttrNVVM shuffle kind
return_value_and_is_valid::mlir::UnitAttrunit attribute

Operands: 

OperandDescription
dst32-bit signless integer
valLLVM dialect-compatible type
offset32-bit signless integer
mask_and_clamp32-bit signless integer

Results: 

ResultDescription
resLLVM dialect-compatible type

nvvm.read.ptx.sreg.tid.x (::mlir::NVVM::ThreadIdXOp) 

Syntax:

operation ::= `nvvm.read.ptx.sreg.tid.x` attr-dict `:` type($res)

Interfaces: NoSideEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Results: 

ResultDescription
resLLVM dialect-compatible type

nvvm.read.ptx.sreg.tid.y (::mlir::NVVM::ThreadIdYOp) 

Syntax:

operation ::= `nvvm.read.ptx.sreg.tid.y` attr-dict `:` type($res)

Interfaces: NoSideEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Results: 

ResultDescription
resLLVM dialect-compatible type

nvvm.read.ptx.sreg.tid.z (::mlir::NVVM::ThreadIdZOp) 

Syntax:

operation ::= `nvvm.read.ptx.sreg.tid.z` attr-dict `:` type($res)

Interfaces: NoSideEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Results: 

ResultDescription
resLLVM dialect-compatible type

nvvm.vote.ballot.sync (::mlir::NVVM::VoteBallotOp) 

Operands: 

OperandDescription
maskLLVM dialect-compatible type
predLLVM dialect-compatible type

Results: 

ResultDescription
resLLVM dialect-compatible type

nvvm.wmma.load (::mlir::NVVM::WMMALoadOp) 

Warp synchronous matrix load

Syntax:

operation ::= `nvvm.wmma.load` $ptr `,` $stride attr-dict `:` functional-type($ptr, $res)

Attributes: 

AttributeMLIR TypeDescription
m::mlir::IntegerAttr32-bit signless integer attribute
n::mlir::IntegerAttr32-bit signless integer attribute
k::mlir::IntegerAttr32-bit signless integer attribute
layoutmlir::StringAttrNVVM MMA layout
eltypemlir::StringAttrNVVM MMA types
fragmlir::StringAttrNVVM MMA frag type

Operands: 

OperandDescription
ptrLLVM pointer type
stride32-bit signless integer

Results: 

ResultDescription
resLLVM structure type

nvvm.wmma.mma (::mlir::NVVM::WMMAMmaOp) 

Warp synchronous matrix-multiply accumulate using tensor cores.

Syntax:

operation ::= `nvvm.wmma.mma` $args attr-dict `:` functional-type($args, $res)

Attributes: 

AttributeMLIR TypeDescription
m::mlir::IntegerAttr32-bit signless integer attribute
n::mlir::IntegerAttr32-bit signless integer attribute
k::mlir::IntegerAttr32-bit signless integer attribute
layoutAmlir::StringAttrNVVM MMA layout
layoutBmlir::StringAttrNVVM MMA layout
eltypeAmlir::StringAttrNVVM MMA types
eltypeBmlir::StringAttrNVVM MMA types

Operands: 

OperandDescription
argsLLVM dialect-compatible type

Results: 

ResultDescription
resLLVM structure type

nvvm.wmma.store (::mlir::NVVM::WMMAStoreOp) 

Warp synchronous matrix store

Syntax:

operation ::= `nvvm.wmma.store` $ptr `,` $stride `,` $args attr-dict `:` type($ptr) `,` type($args)

Attributes: 

AttributeMLIR TypeDescription
m::mlir::IntegerAttr32-bit signless integer attribute
n::mlir::IntegerAttr32-bit signless integer attribute
k::mlir::IntegerAttr32-bit signless integer attribute
layoutmlir::StringAttrNVVM MMA layout
eltypemlir::StringAttrNVVM MMA types

Operands: 

OperandDescription
ptrLLVM pointer type
argsLLVM dialect-compatible type
stride32-bit signless integer

nvvm.read.ptx.sreg.warpsize (::mlir::NVVM::WarpSizeOp) 

Syntax:

operation ::= `nvvm.read.ptx.sreg.warpsize` attr-dict `:` type($res)

Interfaces: NoSideEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Results: 

ResultDescription
resLLVM dialect-compatible type