wmma GPU Virtual ISA NVIDIA

Warp Matrix Multiply-Accumulate Matrix Operations

wmma.load.a.sync.aligned.layout.shape.type r, [p];

Higher-level warp matrix-multiply-accumulate built from explicit load/mma/store steps.

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 6.0
Minimum Target sm_70

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
wmma.load.a.sync.aligned.layout.shape.type r, [p]; sm_70 Load a matrix fragment cooperatively across the warp.
wmma.mma.sync.aligned.layout.shape.dtype.ctype d, a, b, c; sm_70 Perform the accumulate step on already-loaded fragments.
wmma.store.d.sync.aligned.layout.shape.type [p], r; sm_70 Store an accumulator fragment cooperatively across the warp.

Operands

  • r/d
    Fragment register(s)
  • a
    Matrix A fragment
  • b
    Matrix B fragment
  • c
    Accumulator fragment (input)
  • p
    Memory address for load/store forms

At a Glance

Data Types —

Related AMDGPU Concepts

Reference

NVIDIA PTX ISA

Semantics

Together, the load/mma/store triplet compute D = A * B + C for a fixed tile shape, cooperatively across the warp.

Sources