tcgen05.mma GPU Virtual ISA NVIDIA

tcgen05.mma TensorCore 5th Generation of MMA Instructions

// 1. Floating-point type without block scaling:
tcgen05.mma.cta_group.kind [d-tmem], a-desc, b-desc, idesc,
{ disable-output-lane }, enable-input-d {, scale-input-d};

Instruction tcgen05.mma is an asynchronous instruction which initiates an 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.6
Minimum Target sm_100a

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
// 1. Floating-point type without block scaling: tcgen05.mma.cta_group.kind [d-tmem], a-desc, b-desc, idesc, { disable-output-lane }, enable-input-d {, scale-input-d}; sm_100a Instruction tcgen05.mma is an asynchronous instruction which initiates an 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… (see the official PTX ISA docs for the full description)

Operands

At a Glance

Data Types -

Related

More in TensorCore 5th Generation of MMA Instructions

Reference

NVIDIA PTX ISA

Description

Instruction tcgen05.mma is an asynchronous instruction which initiates an 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 operation of the form D = A*B is issued when the input predicate argument enable-input-d is false. The optional immediate argument scale-input-d can be specified to scale the input matrix D as follows: D = A*B+D * (2 ^ - scale-input-d) The valid range of values for argument scale-input-d is [0, 15]. (see the official PTX ISA docs for the full description)

Examples

tcgen05.mma.cta_group::1.kind::tf32      [taddr0],  adesc,  bdesc, idesc, {m0, m1, m2, m3}, p;
tcgen05.mma.cta_group::1.kind::mxf8f6f4  [taddr2],  [taddr1],  bdesc, idesc,
                                         [tmem_scaleA], [tmem_scaleB], p;

tcgen05.commit.cta_group::1.mbarrier::arrive::one.b64 [mbarObj0];
// (truncated - see the official PTX ISA docs for the full example)

Reproduced from NVIDIA's official PTX ISA documentation for technical accuracy.

Sources