v_mfma_f32_16x16x16f16 GPU Native ISA AMD Matrix

V MFMA F32 16X16X16F16 Matrix Core Operations

v_mfma_f32_16x16x16f16 D, A, B, C

Matrix-fused-multiply-add: cooperative 16x16x16 matrix-multiply-accumulate on matrix-core hardware, fp16 inputs, fp32 accumulate.

Encoding

Verified bit-level encoding data is not yet available for this instruction. The instruction-format classification below (VOP3P-MAI) is well-documented and stable; exact per-target opcode/field bit positions have not yet been imported from a verified source.
Format VOP3P-MAI
Width 32 bits
Execution Unit Matrix Core Unit
Register Classes VGPR, AGPR

Operands

  • D
    Accumulator fragment (destination, VGPR range)
  • A
    Matrix A fragment (VGPR range)
  • B
    Matrix B fragment (VGPR range)
  • C
    Accumulator fragment (input, VGPR range)

GFX Target Compatibility

TargetSupport
gfx942✅ Supported
gfx1100— Not implemented

Matrix-core (MFMA) instructions are a CDNA-family feature; gfx1100 (RDNA3) does not implement this specific MFMA variant — see the AMDGPU landing page's target compatibility notes.

Related PTX Concepts

Reference

AMDGPU / GFX ISA

Semantics

D = A * B + C for a 16x16x16 tile, computed cooperatively across the wavefront; operand fragments are distributed across lanes per a hardware-defined layout specific to the MFMA family and target.

Sources