'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)
Traits: AlwaysSpeculatableImplTrait
Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Results: ¶
Result | Description |
---|---|
res | LLVM dialect-compatible type |
nvvm.read.ptx.sreg.ntid.y
(::mlir::NVVM::BlockDimYOp) ¶
Syntax:
operation ::= `nvvm.read.ptx.sreg.ntid.y` attr-dict `:` type($res)
Traits: AlwaysSpeculatableImplTrait
Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Results: ¶
Result | Description |
---|---|
res | LLVM dialect-compatible type |
nvvm.read.ptx.sreg.ntid.z
(::mlir::NVVM::BlockDimZOp) ¶
Syntax:
operation ::= `nvvm.read.ptx.sreg.ntid.z` attr-dict `:` type($res)
Traits: AlwaysSpeculatableImplTrait
Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Results: ¶
Result | Description |
---|---|
res | LLVM dialect-compatible type |
nvvm.read.ptx.sreg.ctaid.x
(::mlir::NVVM::BlockIdXOp) ¶
Syntax:
operation ::= `nvvm.read.ptx.sreg.ctaid.x` attr-dict `:` type($res)
Traits: AlwaysSpeculatableImplTrait
Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Results: ¶
Result | Description |
---|---|
res | LLVM dialect-compatible type |
nvvm.read.ptx.sreg.ctaid.y
(::mlir::NVVM::BlockIdYOp) ¶
Syntax:
operation ::= `nvvm.read.ptx.sreg.ctaid.y` attr-dict `:` type($res)
Traits: AlwaysSpeculatableImplTrait
Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Results: ¶
Result | Description |
---|---|
res | LLVM dialect-compatible type |
nvvm.read.ptx.sreg.ctaid.z
(::mlir::NVVM::BlockIdZOp) ¶
Syntax:
operation ::= `nvvm.read.ptx.sreg.ctaid.z` attr-dict `:` type($res)
Traits: AlwaysSpeculatableImplTrait
Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Results: ¶
Result | Description |
---|---|
res | LLVM dialect-compatible type |
nvvm.cp.async.commit.group
(::mlir::NVVM::CpAsyncCommitGroupOp) ¶
Syntax:
operation ::= `nvvm.cp.async.commit.group` attr-dict
nvvm.cp.async.shared.global
(::mlir::NVVM::CpAsyncOp) ¶
Syntax:
operation ::= `nvvm.cp.async.shared.global` $dst `,` $src `,` $size attr-dict `:` type(operands)
Attributes: ¶
Attribute | MLIR Type | Description |
---|---|---|
size | ::mlir::IntegerAttr | 32-bit signless integer attribute |
bypass_l1 | ::mlir::UnitAttr | unit attribute |
Operands: ¶
Operand | Description |
---|---|
dst | LLVM pointer to 8-bit signless integer |
src | LLVM pointer to 8-bit signless integer |
nvvm.cp.async.wait.group
(::mlir::NVVM::CpAsyncWaitGroupOp) ¶
Syntax:
operation ::= `nvvm.cp.async.wait.group` $n attr-dict
Attributes: ¶
Attribute | MLIR Type | Description |
---|---|---|
n | ::mlir::IntegerAttr | 32-bit signless integer attribute |
nvvm.read.ptx.sreg.nctaid.x
(::mlir::NVVM::GridDimXOp) ¶
Syntax:
operation ::= `nvvm.read.ptx.sreg.nctaid.x` attr-dict `:` type($res)
Traits: AlwaysSpeculatableImplTrait
Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Results: ¶
Result | Description |
---|---|
res | LLVM dialect-compatible type |
nvvm.read.ptx.sreg.nctaid.y
(::mlir::NVVM::GridDimYOp) ¶
Syntax:
operation ::= `nvvm.read.ptx.sreg.nctaid.y` attr-dict `:` type($res)
Traits: AlwaysSpeculatableImplTrait
Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Results: ¶
Result | Description |
---|---|
res | LLVM dialect-compatible type |
nvvm.read.ptx.sreg.nctaid.z
(::mlir::NVVM::GridDimZOp) ¶
Syntax:
operation ::= `nvvm.read.ptx.sreg.nctaid.z` attr-dict `:` type($res)
Traits: AlwaysSpeculatableImplTrait
Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Results: ¶
Result | Description |
---|---|
res | LLVM dialect-compatible type |
nvvm.read.ptx.sreg.laneid
(::mlir::NVVM::LaneIdOp) ¶
Syntax:
operation ::= `nvvm.read.ptx.sreg.laneid` attr-dict `:` type($res)
Traits: AlwaysSpeculatableImplTrait
Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Results: ¶
Result | Description |
---|---|
res | LLVM dialect-compatible type |
nvvm.ldmatrix
(::mlir::NVVM::LdMatrixOp) ¶
cooperative matrix load
Syntax:
operation ::= `nvvm.ldmatrix` $ptr attr-dict `:` functional-type($ptr, $res)
Attributes: ¶
Attribute | MLIR Type | Description |
---|---|---|
num | ::mlir::IntegerAttr | 32-bit signless integer attribute |
layout | ::mlir::NVVM::MMALayoutAttr | NVVM MMA layout |
Operands: ¶
Operand | Description |
---|---|
ptr | LLVM pointer type |
Results: ¶
Result | Description |
---|---|
res | any type |
nvvm.mma.sync
(::mlir::NVVM::MmaOp) ¶
cooperative matrix-multiply and accumulate
The nvvm.mma.sync
operation collectively performs the operation
D = matmul(A, B) + C
using all threads in a warp.
All the threads in the warp must execute the same mma.sync
operation.
For each possible multiplicand PTX data type, there are one or more possible instruction shapes given as “mMnNkK”. The below table describes the posssibilities as well as the types required for the operands. Note that the data type for C (the accumulator) and D (the result) can vary independently when there are multiple possibilities in the “C/D Type” column.
When an optional attribute cannot be immediately inferred from the types of the operands and the result during parsing or validation, an error will be raised.
b1Op
is only relevant when the binary (b1) type is given to
multiplicandDataType
. It specifies how the multiply-and-acumulate is
performed and is either xor_popc
or and_poc
. The default is xor_popc
.
intOverflowBehavior
is only relevant when the multiplicandType
attribute
is one of u8, s8, u4, s4
, this attribute describes how overflow is handled
in the accumulator. When the attribute is satfinite
, the accumulator values
are clamped in the int32 range on overflow. This is the default behavior.
Alternatively, accumulator behavior wrapped
can also be specified, in
which case overflow wraps from one end of the range to the other.
layoutA
and layoutB
are required and should generally be set to
#nvvm.mma_layout<row>
and #nvvm.mma_layout<col>
respectively, but other
combinations are possible for certain layouts according to the table below.
| A/B Type | Shape | ALayout | BLayout | A Type | B Type | C/D Type |
|----------|-----------|---------|---------|----------|----------|-------------------|
| f64 | .m8n8k4 | row | col | 1x f64 | 1x f64 | 2x f64 |
| f16 | .m8n8k4 | row/col | row/col | 2x f16x2 | 2x f16x2 | 4x f16x2 or 8xf32 |
| | .m16n8k8 | row | col | 2x f16x2 | 1x f16x2 | 2x f16x2 or 4 f32 |
| | .m16n8k16 | row | col | 4x f16x2 | 2x f16x2 | 2x f16x2 or 4 f32 |
| bf16 | .m16n8k8 | row | col | 2x f16x2 | 1x f16x2 | 2x f16x2 or 4 f32 |
| | .m16n8k16 | row | col | 4x f16x2 | 2x f16x2 | 2x f16x2 or 4 f32 |
| tf32 | .m16n8k4 | row | col | 2x i32 | 1x i32 | 4x f32 |
| | .m16n8k8 | row | col | 4x i32 | 2x i32 | 2x f16x2 or 4 f32 |
| u8/s8 | .m8n8k16 | row | col | 1x i32 | 1x i32 | 2x i32 |
| | .m16n8k16 | row | col | 2x i32 | 1x i32 | 4x i32 |
| | .m16n8k32 | row | col | 4x i32 | 2x i32 | 4x i32 |
| u4/s4 | .m8n8k32 | row | col | 1x i32 | 1x i32 | 2x i32 |
| | m16n8k32 | row | col | 2x i32 | 1x i32 | 4x i32 |
| | m16n8k64 | row | col | 4x i32 | 2x i32 | 4x i32 |
| b1 | m8n8k128 | row | col | 1x i32 | 1x i32 | 2x i32 |
| | m16n8k128 | row | col | 2x i32 | 1x i32 | 4x i32 |
Example:
%128 = nvvm.mma.sync A[%120, %121, %122, %123]
B[%124, %125]
C[%126, %127]
{layoutA = #nvvm.mma_layout<row>,
layoutB = #nvvm.mma_layout<col>,
shape = {k = 16 : i32, m = 16 : i32, n = 8 : i32}}
: (vector<2xf16>, vector<2xf16>, vector<2xf16>)
-> !llvm.struct<(vector<2xf16>, vector<2xf16>)>
Traits: AttrSizedOperandSegments
Attributes: ¶
Attribute | MLIR Type | Description |
---|---|---|
shape | ::mlir::NVVM::MMAShapeAttr | Attribute for MMA operation shape. |
b1Op | ::mlir::NVVM::MMAB1OpAttr | MMA binary operations |
intOverflowBehavior | ::mlir::NVVM::MMAIntOverflowAttr | MMA overflow options |
layoutA | ::mlir::NVVM::MMALayoutAttr | NVVM MMA layout |
layoutB | ::mlir::NVVM::MMALayoutAttr | NVVM MMA layout |
multiplicandAPtxType | ::mlir::NVVM::MMATypesAttr | NVVM MMA types |
multiplicandBPtxType | ::mlir::NVVM::MMATypesAttr | NVVM MMA types |
Operands: ¶
Operand | Description |
---|---|
operandA | LLVM dialect-compatible type |
operandB | LLVM dialect-compatible type |
operandC | LLVM dialect-compatible type |
Results: ¶
Result | Description |
---|---|
res | LLVM structure type |
nvvm.rcp.approx.ftz.f
(::mlir::NVVM::RcpApproxFtzF32Op) ¶
Syntax:
operation ::= `nvvm.rcp.approx.ftz.f` $arg attr-dict `:` type($res)
Traits: AlwaysSpeculatableImplTrait
Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands: ¶
Operand | Description |
---|---|
arg | 32-bit float |
Results: ¶
Result | Description |
---|---|
res | 32-bit float |
nvvm.redux.sync
(::mlir::NVVM::ReduxOp) ¶
Syntax:
operation ::= `nvvm.redux.sync` $kind $val `,` $mask_and_clamp attr-dict `:` type($val) `->` type($res)
Attributes: ¶
Attribute | MLIR Type | Description |
---|---|---|
kind | ::mlir::NVVM::ReduxKindAttr | NVVM redux kind |
Operands: ¶
Operand | Description |
---|---|
val | LLVM dialect-compatible type |
mask_and_clamp | 32-bit signless integer |
Results: ¶
Result | Description |
---|---|
res | LLVM 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: ¶
Attribute | MLIR Type | Description |
---|---|---|
kind | ::mlir::NVVM::ShflKindAttr | NVVM shuffle kind |
return_value_and_is_valid | ::mlir::UnitAttr | unit attribute |
Operands: ¶
Operand | Description |
---|---|
dst | 32-bit signless integer |
val | LLVM dialect-compatible type |
offset | 32-bit signless integer |
mask_and_clamp | 32-bit signless integer |
Results: ¶
Result | Description |
---|---|
res | LLVM dialect-compatible type |
nvvm.bar.warp.sync
(::mlir::NVVM::SyncWarpOp) ¶
Syntax:
operation ::= `nvvm.bar.warp.sync` $mask attr-dict `:` type($mask)
Operands: ¶
Operand | Description |
---|---|
mask | LLVM dialect-compatible type |
nvvm.read.ptx.sreg.tid.x
(::mlir::NVVM::ThreadIdXOp) ¶
Syntax:
operation ::= `nvvm.read.ptx.sreg.tid.x` attr-dict `:` type($res)
Traits: AlwaysSpeculatableImplTrait
Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Results: ¶
Result | Description |
---|---|
res | LLVM dialect-compatible type |
nvvm.read.ptx.sreg.tid.y
(::mlir::NVVM::ThreadIdYOp) ¶
Syntax:
operation ::= `nvvm.read.ptx.sreg.tid.y` attr-dict `:` type($res)
Traits: AlwaysSpeculatableImplTrait
Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Results: ¶
Result | Description |
---|---|
res | LLVM dialect-compatible type |
nvvm.read.ptx.sreg.tid.z
(::mlir::NVVM::ThreadIdZOp) ¶
Syntax:
operation ::= `nvvm.read.ptx.sreg.tid.z` attr-dict `:` type($res)
Traits: AlwaysSpeculatableImplTrait
Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Results: ¶
Result | Description |
---|---|
res | LLVM dialect-compatible type |
nvvm.vote.ballot.sync
(::mlir::NVVM::VoteBallotOp) ¶
Operands: ¶
Operand | Description |
---|---|
mask | LLVM dialect-compatible type |
pred | LLVM dialect-compatible type |
Results: ¶
Result | Description |
---|---|
res | LLVM 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: ¶
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 |
layout | ::mlir::NVVM::MMALayoutAttr | NVVM MMA layout |
eltype | ::mlir::NVVM::MMATypesAttr | NVVM MMA types |
frag | ::mlir::NVVM::MMAFragAttr | NVVM MMA frag type |
Operands: ¶
Operand | Description |
---|---|
ptr | LLVM pointer type |
stride | 32-bit signless integer |
Results: ¶
Result | Description |
---|---|
res | LLVM 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: ¶
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 |
layoutA | ::mlir::NVVM::MMALayoutAttr | NVVM MMA layout |
layoutB | ::mlir::NVVM::MMALayoutAttr | NVVM MMA layout |
eltypeA | ::mlir::NVVM::MMATypesAttr | NVVM MMA types |
eltypeB | ::mlir::NVVM::MMATypesAttr | NVVM MMA types |
Operands: ¶
Operand | Description |
---|---|
args | LLVM dialect-compatible type |
Results: ¶
Result | Description |
---|---|
res | LLVM structure type |
nvvm.wmma.store
(::mlir::NVVM::WMMAStoreOp) ¶
Warp synchronous matrix store
Syntax:
operation ::= `nvvm.wmma.store` $ptr `,` $stride `,` $args attr-dict `:` qualified(type($ptr)) `,`
type($args)
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 |
layout | ::mlir::NVVM::MMALayoutAttr | NVVM MMA layout |
eltype | ::mlir::NVVM::MMATypesAttr | NVVM MMA types |
Operands: ¶
Operand | Description |
---|---|
ptr | LLVM pointer type |
args | LLVM dialect-compatible type |
stride | 32-bit signless integer |
nvvm.read.ptx.sreg.warpsize
(::mlir::NVVM::WarpSizeOp) ¶
Syntax:
operation ::= `nvvm.read.ptx.sreg.warpsize` attr-dict `:` type($res)
Traits: AlwaysSpeculatableImplTrait
Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Results: ¶
Result | Description |
---|---|
res | LLVM dialect-compatible type |
Attribute definition ¶
MMAB1OpAttr ¶
MMA binary operations
Syntax:
#nvvm.mma_b1op<
::mlir::NVVM::MMAB1Op # value
>
Parameters: ¶
Parameter | C++ type | Description |
---|---|---|
value | ::mlir::NVVM::MMAB1Op | an enum of type MMAB1Op |
MMAFragAttr ¶
NVVM MMA frag type
Syntax:
#nvvm.mma_frag<
::mlir::NVVM::MMAFrag # value
>
Parameters: ¶
Parameter | C++ type | Description |
---|---|---|
value | ::mlir::NVVM::MMAFrag | an enum of type MMAFrag |
MMAIntOverflowAttr ¶
MMA overflow options
Syntax:
#nvvm.mma_int_overflow<
::mlir::NVVM::MMAIntOverflow # value
>
Parameters: ¶
Parameter | C++ type | Description |
---|---|---|
value | ::mlir::NVVM::MMAIntOverflow | an enum of type MMAIntOverflow |
MMALayoutAttr ¶
NVVM MMA layout
Syntax:
#nvvm.mma_layout<
::mlir::NVVM::MMALayout # value
>
Parameters: ¶
Parameter | C++ type | Description |
---|---|---|
value | ::mlir::NVVM::MMALayout | an enum of type MMALayout |
MMATypesAttr ¶
NVVM MMA types
Syntax:
#nvvm.mma_type<
::mlir::NVVM::MMATypes # value
>
Parameters: ¶
Parameter | C++ type | Description |
---|---|---|
value | ::mlir::NVVM::MMATypes | an enum of type MMATypes |
MMAShapeAttr ¶
Attribute for MMA operation shape.
Syntax:
#nvvm.shape<
int, # m
int, # n
int # k
>
Parameters: ¶
Parameter | C++ type | Description |
---|---|---|
m | int | |
n | int | |
k | int |
ReduxKindAttr ¶
NVVM redux kind
Syntax:
#nvvm.redux_kind<
::mlir::NVVM::ReduxKind # value
>
Parameters: ¶
Parameter | C++ type | Description |
---|---|---|
value | ::mlir::NVVM::ReduxKind | an enum of type ReduxKind |
ShflKindAttr ¶
NVVM shuffle kind
Syntax:
#nvvm.shfl_kind<
::mlir::NVVM::ShflKind # value
>
Parameters: ¶
Parameter | C++ type | Description |
---|---|---|
value | ::mlir::NVVM::ShflKind | an enum of type ShflKind |