'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: ¶

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

Syntax:

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

Results: ¶

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

Syntax:

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

Results: ¶

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

Syntax:

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

Results: ¶

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

Syntax:

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

Results: ¶

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

Syntax:

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

Results: ¶

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

Syntax:

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

Results: ¶

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

Syntax:

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

Results: ¶

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

Syntax:

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

Results: ¶

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

Syntax:

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

Results: ¶

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

Syntax:

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

Operands: ¶

Results: ¶

nvvm.shfl.sync.bfly (::mlir::NVVM::ShflBflyOp) ¶

Attributes: ¶

Operands: ¶

Results: ¶

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

Syntax:

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

Results: ¶

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

Syntax:

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

Results: ¶

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

Syntax:

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

Results: ¶

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

Operands: ¶

Results: ¶

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: ¶

Results: ¶

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: ¶

Results: ¶

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: ¶

Results: ¶

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: ¶

Results: ¶

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: ¶

Results: ¶

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: ¶

Results: ¶

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: ¶

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: ¶

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

Syntax:

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

Results: ¶