Matrix Multiply (Tensor/Matrix-Core) Matrix Operations

Compute a small matrix-multiply-accumulate using dedicated tensor/matrix-core hardware, cooperatively across a warp/wavefront.

Vendor-Neutral Definition

D = A * B + C for fixed tile shapes (e.g. 16x16x16), where operand fragments are distributed across the lanes of a warp/wavefront per a hardware-defined layout.
⚠️ Cross-vendor mappings describe semantic relationships and are not guaranteed one-to-one compiler translations.
unverified Both perform a cooperative, hardware-accelerated matrix-multiply-accumulate across a warp/wavefront, but shape and layout parity has not been individually verified.
Property NVIDIA PTX AMDGPU
Instruction mma, wmma v_mfma_f32_16x16x16f16
ISA Layer Virtual Native
Data Types — f32
Version / Target Introduced PTX ISA 6.4, PTX ISA 6.0 gfx942

Important Differences

  • Tile shapes, operand-fragment-to-lane layout, and supported input/accumulator type combinations differ substantially between NVIDIA's mma/wmma family and AMD's MFMA family and are generation-specific on both sides.
  • Treat this as a conceptual pointer only — do not assume a specific PTX mma shape maps 1:1 onto v_mfma_f32_16x16x16f16 without checking both vendors' current documentation for that exact shape.

Source Evidence

Verification method: documentation