mma GPU Virtual ISA NVIDIA

Matrix Multiply-Accumulate (Tensor Core) Matrix Operations

mma.sync.aligned.shape.row.col.dtype.atype.btype.ctype d, a, b, c;

Cooperative, warp-wide matrix-multiply-accumulate executed on tensor-core hardware.

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.4
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
mma.sync.aligned.shape.row.col.dtype.atype.btype.ctype d, a, b, c; sm_70 Synchronizing warp-wide MMA for a fixed tile shape (e.g. m16n8k16); operand fragments are distributed across the warp's lanes per a hardware-defined layout.

Operands

  • d
    Accumulator fragment (destination)
  • a
    Matrix A fragment
  • b
    Matrix B fragment
  • c
    Accumulator fragment (input)

At a Glance

Data Types —

Related AMDGPU Concepts

Reference

NVIDIA PTX ISA

Semantics

D = A * B + C for the fixed tile shape, computed cooperatively across all 32 lanes of the warp.

Sources