wgmma.mma_async.sp GPU Virtual ISA NVIDIA

wgmma.mma_async.sp Asynchronous Warpgroup Level Matrix Multiply-Accumulate Instructions

wgmma.mma_async.sp.sync.aligned.shape.dtype.f16.f16 d, a-desc, b-desc, sp-meta, sp-sel, scale-d, imm-scale-a, imm-scale-b, imm-trans-a, imm-trans-b;

Instruction wgmma.mma_async issues a MxNxK matrix multiply and accumulate operation, D = A*B+D, where the A matrix is MxK, the B matrix is KxN, and the D matrix is MxN.

Encoding

PTX is a virtual instruction set. It has no single, stable native binary encoding - the compiler lowers this instruction to different native machine code depending on the selected NVIDIA target architecture (compute capability). This page intentionally shows no bit-diagram; see the target/version requirements below for what governs how this instruction compiles.
PTX ISA Version Introduced PTX ISA 8.2
Minimum Target sm_90a

Syntax Forms

One mnemonic covers many type / state-space / scope / modifier combinations - each row below is an independently valid form.

Syntax Data Types State Space(s) Modifiers Min. Target Description
wgmma.mma_async.sp.sync.aligned.shape.dtype.f16.f16 d, a-desc, b-desc, sp-meta, sp-sel, scale-d, imm-scale-a, imm-scale-b, imm-trans-a, imm-trans-b; sm_90a Instruction wgmma.mma_async issues a MxNxK matrix multiply and accumulate operation, D = A*B+D, where the A matrix is MxK, the B matrix is KxN, and the D matrix is MxN. (see the official PTX ISA docs for the full description)

Operands

  • d
    Destination register
  • a-desc
    Operand
  • b-desc
    Operand
  • sp-meta
    Operand
  • sp-sel
    Operand
  • scale-d
    Operand
  • imm-scale-a
    Operand
  • imm-scale-b
    Operand
  • imm-trans-a
    Operand
  • imm-trans-b
    Operand

At a Glance

Data Types -

Related

More in Asynchronous Warpgroup Level Matrix Multiply-Accumulate Instructions

Reference

NVIDIA PTX ISA

Description

Instruction wgmma.mma_async issues a MxNxK matrix multiply and accumulate operation, D = A*B+D, where the A matrix is MxK, the B matrix is KxN, and the D matrix is MxN. The matrix A is stored in the packed format Mx(K/2) as described in Sparse matrix storage. The operation of the form D = A*B is issued when the input predicate argument scale-d is false. wgmma.fence instruction must be used to fence the register accesses of wgmma.mma_async instruction from their prior accesses. (see the official PTX ISA docs for the full description)

Sources