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)

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)

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)

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)

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)

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)

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)

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)

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)

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)

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.bfly (::mlir::NVVM::ShflBflyOp) 

Attributes: 

AttributeMLIR TypeDescription
return_value_and_is_valid::mlir::UnitAttrunit attribute

Operands: 

OperandDescription
dstLLVM dialect-compatible type
valLLVM dialect-compatible type
offsetLLVM dialect-compatible type
mask_and_clampLLVM dialect-compatible type

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)

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)

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)

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.m16n16k16.load.a.f16.row.stride (::mlir::NVVM::WMMALoadAM16N16K16Op) 

Warp synchronous matrix load

Syntax:

operation ::= `nvvm.wmma.m16n16k16.load.a.f16.row.stride` $args attr-dict `:` functional-type($args, $res)

“The nvvm.wmma.m*n*k*.load.[a, b, c] operation” “loads a matrix collectively using all the threads in a warp.”

"The operation takes two arguments, the address from where the matrix"
"elements are to be loaded from and a stride. The stride argument"
"represents the leading dimension of the source matrix. The address and"
"the stride are required to be the same across all threads in the warp."
"Each thread in a warp holds a certain number of elements. The Op returns"
"a LLVMStruct which holds the elements of the matrix held by this thread."

"This op is meant to be used along with `nvvm.wmma.m*n*k*.store` and"
"`nvvm.wmma.m*n*k*.mma`."
Example:

```mlir
%2 = nvvm.wmma.m16n16k16.load.a %0, %1 : !llvm.ptr<i32, 3>, !llvm.i32 ->
!llvm.struct<(vec<2 x half>, vec<2 x half>, vec<2 x half>, vec<2 x half>,
vec<2 x half>, vec<2 x half>, vec<2 x half>, vec<2 x half>)>
```

Operands: 

OperandDescription
argsLLVM dialect-compatible type

Results: 

ResultDescription
resLLVM structure type

nvvm.wmma.m16n16k16.load.b.f16.row.stride (::mlir::NVVM::WMMALoadBM16N16K16Op) 

Warp synchronous matrix load

Syntax:

operation ::= `nvvm.wmma.m16n16k16.load.b.f16.row.stride` $args attr-dict `:` functional-type($args, $res)

“The nvvm.wmma.m*n*k*.load.[a, b, c] operation” “loads a matrix collectively using all the threads in a warp.”

"The operation takes two arguments, the address from where the matrix"
"elements are to be loaded from and a stride. The stride argument"
"represents the leading dimension of the source matrix. The address and"
"the stride are required to be the same across all threads in the warp."
"Each thread in a warp holds a certain number of elements. The Op returns"
"a LLVMStruct which holds the elements of the matrix held by this thread."

"This op is meant to be used along with `nvvm.wmma.m*n*k*.store` and"
"`nvvm.wmma.m*n*k*.mma`."
Example:

```mlir
%2 = nvvm.wmma.m16n16k16.load.b %0, %1 : !llvm.ptr<i32, 3>, !llvm.i32 ->
!llvm.struct<(vec<2 x half>, vec<2 x half>, vec<2 x half>, vec<2 x half>,
vec<2 x half>, vec<2 x half>, vec<2 x half>, vec<2 x half>)>
```

Operands: 

OperandDescription
argsLLVM dialect-compatible type

Results: 

ResultDescription
resLLVM structure type

nvvm.wmma.m16n16k16.load.c.f16.row.stride (::mlir::NVVM::WMMALoadCF16M16N16K16Op) 

Warp synchronous matrix load

Syntax:

operation ::= `nvvm.wmma.m16n16k16.load.c.f16.row.stride` $args attr-dict `:` functional-type($args, $res)

“The nvvm.wmma.m*n*k*.load.[a, b, c] operation” “loads a matrix collectively using all the threads in a warp.”

"The operation takes two arguments, the address from where the matrix"
"elements are to be loaded from and a stride. The stride argument"
"represents the leading dimension of the source matrix. The address and"
"the stride are required to be the same across all threads in the warp."
"Each thread in a warp holds a certain number of elements. The Op returns"
"a LLVMStruct which holds the elements of the matrix held by this thread."

"This op is meant to be used along with `nvvm.wmma.m*n*k*.store` and"
"`nvvm.wmma.m*n*k*.mma`."
Example:

```mlir
%2 = nvvm.wmma.m16n16k16.load.c.f16.row.stride %0, %1 : !llvm.ptr<i32, 3>, !llvm.i32 ->
!llvm.struct<(vec<2 x half>, vec<2 x half>, vec<2 x half>, vec<2 x half>)>
```

Operands: 

OperandDescription
argsLLVM dialect-compatible type

Results: 

ResultDescription
resLLVM structure type

nvvm.wmma.m16n16k16.load.c.f32.row.stride (::mlir::NVVM::WMMALoadCF32M16N16K16Op) 

Warp synchronous matrix load

Syntax:

operation ::= `nvvm.wmma.m16n16k16.load.c.f32.row.stride` $args attr-dict `:` functional-type($args, $res)

“The nvvm.wmma.m*n*k*.load.[a, b, c] operation” “loads a matrix collectively using all the threads in a warp.”

"The operation takes two arguments, the address from where the matrix"
"elements are to be loaded from and a stride. The stride argument"
"represents the leading dimension of the source matrix. The address and"
"the stride are required to be the same across all threads in the warp."
"Each thread in a warp holds a certain number of elements. The Op returns"
"a LLVMStruct which holds the elements of the matrix held by this thread."

"This op is meant to be used along with `nvvm.wmma.m*n*k*.store` and"
"`nvvm.wmma.m*n*k*.mma`."
Example:

```mlir
%2 = nvvm.wmma.m16n16k16.load.c.f32.row.stride %0, %1 : !llvm.ptr<i32, 3>, !llvm.i32 ->
!llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)>
```

Operands: 

OperandDescription
argsLLVM dialect-compatible type

Results: 

ResultDescription
resLLVM structure type

nvvm.wmma.m16n16k16.mma.row.row.f16.f16 (::mlir::NVVM::WMMAMmaF16F16M16N16K16Op) 

Warp synchronous matrix-multiply accumulate using tensor cores.

The nvvm.wmma.m*n*k*.mma operation performs a matrix-multiply accumulate (mma) operation using all the threads in a warp.

The operation performed is represented as D = A * B + C. The operation takes as arguments the elements of the matrices A, B, C and D, held by the current thread. The op returns a LLVM struct which holds a part of the result held by the current thread.

This op is meant to be used along with nvvm.wmma.m16n16k16.load and nvvm.wmma. m16n16k16.store.

Example:

%20 = nvvm.wmma.m16n16k16.mma.row.row.f16.f16 %0, %1, %2, %3, %4, %5, %6, %7, %8,
%9, %10, %11, %12, %13, %14, %15, %16, %17, %18, %19 : vector<2xf16> -> !llvm.struct
<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>

Operands: 

OperandDescription
argsLLVM dialect-compatible type

Results: 

ResultDescription
resLLVM structure type

nvvm.wmma.m16n16k16.mma.row.row.f32.f32 (::mlir::NVVM::WMMAMmaF32F32M16N16K16Op) 

Warp synchronous matrix-multiply accumulate using tensor cores.

Syntax:

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

The nvvm.wmma.m*n*k*.mma operation performs a matrix-multiply accumulate (mma) operation using all the threads in a warp.

The operation performed is represented as D = A * B + C. The operation takes as arguments the elements of the matrices A, B, C and D, held by the current thread. The op returns a LLVM struct which holds a part of the result held by the current thread.

This op is meant to be used along with nvvm.wmma.m16n16k16.load and nvvm.wmma. m16n16k16.store.

Example:

%24 = nvvm.wmma.m16n16k16.mma.row.row.f32.f32 %0, %1, %2, %3, %4, %5, %6, %7, %8
%9, %10, %11, %12, %13, %14, %15, %16, %17, %18, %19, %20, %21, %22, %23 :
(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>,
vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>,
vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>,
vector<2xf16>, f32, f32, f32, f32, f32, f32, f32, f32) -> !llvm.struct<(f32,
f32, f32, f32, f32, f32, f32, f32)>

Operands: 

OperandDescription
argsLLVM dialect-compatible type

Results: 

ResultDescription
resLLVM structure type

nvvm.wmma.m16n16k16.store.d.f16.row.stride (::mlir::NVVM::WMMAStoreF16M16N16K16Op) 

Warp synchronous matrix store

Syntax:

operation ::= `nvvm.wmma.m16n16k16.store.d.f16.row.stride` $args attr-dict `:` type($args)

The nvvm.wmma.m*n*k*.store operation stores a matrix collectively using all the threads in a warp.

The operation takes as arguments the address to where the matrix elements are to be stored, a stride and the elements to store, held by the current thread. The stride argument represents the leading dimension of the destination matrix. The address and the stride are required to be the same across all threads in the warp.

This op is meant to be used along with nvvm.wmma.m16n16k16.load and nvvm.wmma.m16n16k16.mma.

Example:

nvvm.wmma.m16n16k16.stored.f16.row.stride %0, %1, %2, %3, %4, %5, %6 : !llvm.ptr<i32, 3>,
!llvm.struct<(vec<2 x half>, vec<2 x half>, vec<2 x half>, vec<2 x half>)>, !llvm.i32

Operands: 

OperandDescription
argsLLVM dialect-compatible type

nvvm.wmma.m16n16k16.store.d.f32.row.stride (::mlir::NVVM::WMMAStoreF32M16N16K16Op) 

Warp synchronous matrix store

Syntax:

operation ::= `nvvm.wmma.m16n16k16.store.d.f32.row.stride` $args attr-dict `:` type($args)

The nvvm.wmma.m*n*k*.store operation stores a matrix collectively using all the threads in a warp.

The operation takes as arguments the address to where the matrix elements are to be stored, a stride and the elements to store, held by the current thread. The stride argument represents the leading dimension of the destination matrix. The address and the stride are required to be the same across all threads in the warp.

This op is meant to be used along with nvvm.wmma.m16n16k16.load and nvvm.wmma.m16n16k16.mma.

Example:

nvvm.wmma.m16n16k16.store.d.f32.row.stride %0, %1, %2, %3, %4, %5, %6, %7, %8, %9,
%10 : !llvm.ptr<i32, 3>, !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)>,
!llvm.i32

Operands: 

OperandDescription
argsLLVM dialect-compatible type

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

Syntax:

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

Results: 

ResultDescription
resLLVM dialect-compatible type