Skip to main content

CDNA 3 and CDNA 4

Kernel and compiler reference · gfx942 → gfx950

A comparison for GPU kernel and compiler development, with instruction-specific conditions and primary sources.

Research snapshot: November 13, 2026. Independently organized from AMD documentation and LLVM source; not an AMD publication. No GPU benchmark results are included.

Quick comparison

Covered targets: CDNA 3 gfx942 and CDNA 4 gfx950. Start with the four changes below, then follow each link for its scope and evidence.

  • FP8 values need a format-aware migration.

    CDNA 3 uses FNUZ; CDNA 4 uses OCP FP8/BF8. The same byte can mean a different number. Formats and byte example.

  • More work per MFMA changes the input fragment.

    CDNA 4 adds wider-K FP16/BF16 and INT8 instructions and scaled FP8/FP6/FP4 operations. A shared instruction can retain its layout; a new shape needs a new input mapping. Instruction families · Lane formulas.

  • LDS capacity and data paths change together.

    64 KiB / 32 banks becomes 160 KiB / 64 banks per CU. CDNA 4 also adds wider global-to-LDS transfers and LDS reads with transpose. Memory comparison and diagram.

  • Completion, rendezvous and arithmetic hazards remain separate.

    A memory wait, a workgroup barrier and MFMA dependency scheduling solve different problems. Direct-to-LDS copies still need completion tracking. Dependency overview · MFMA compiler cases.

1. Scope and source revisions

This reference follows kernel-visible differences, with product details used only to put the architecture in context. The target scope is deliberately limited to gfx942 and gfx950; it does not cover every gfx9 target, CDNA 5, or RDNA.

Default layout example: wave64, dense 16×16×16 FP16 → FP32 MFMA, one output block, CBSZ=ABID=BLGP=0. The wider-K example changes K to 32 on CDNA 4. Sparse, FP64 and scaled operations have separate layouts and conditions.

Target scope and resource boundaries
PropertyCDNA 3 · gfx942CDNA 4 · gfx950
Processor scopeCDNA 3 · gfx942gfx942; MI300A, MI300X and MI325X are product examples.CDNA 4 · gfx950gfx950; MI350X and MI355X are product examples.
WavefrontCDNA 3 · gfx94264 lanes; the MFMA examples use one full wave.CDNA 4 · gfx95064 lanes; the MFMA examples use one full wave.
LDS per CUCDNA 3 · gfx94264 KiB; 32 banks.

ISA §2.2.1

CDNA 4 · gfx950160 KiB; 64 banks.

ISA §2.2.1

Product contextCDNA 3 · gfx942MI300X: 8 XCDs, 304 active CUs, 4 I/O dies. MI300A uses a different package configuration.

MI300 microarchitecture

CDNA 4 · gfx950MI350X / MI355X: 8 XCDs, 256 active CUs, 2 I/O dies. These counts are not a definition of every CDNA 4 product.

MI350 microarchitecture

The CDNA 3 ISA and CDNA 4 ISA reviewed here both carry 5 August 2025 on their covers. LLVM/Clang source checks use commit 4e9d349ab86d. Web sources were checked on 13 November 2026. PDF links use physical file pages, counted from 1.

Evidence labels and scope

Read each status with its named targets, source revision and instruction scope. These four checks answer separate questions.

ISA inventory

Listed / Not listed reports whether the specified instruction appears in the cited manual or machine-readable inventory. The source is named explicitly, for example AMD XML. An omitted entry does not establish hardware absence.

LLVM target support

Enabled / Not enabled refers to the required builtin or native-instruction features in the exact processor's default feature set at the pinned revision. Forced overrides, software implementations, library dispatch and runtime support are separate questions.

Mechanism

Documented applies only to the behavior described in the cited source. Details unspecified means that source leaves the relevant mechanism or conditions open. Mechanism equivalence unconfirmed means the evidence does not establish that two reported changes share an implementation.

Validation

Source inspected, CPU checked and GPU tested identify the work actually completed, with its stated coverage. Not executed means the named test was not run. Source review, including reading regression expectations, is separate from executing those tests.

Feature reports: Confirmed by AMD means AMD explicitly reports the feature for the named generation or product. Support unconfirmed means the reviewed evidence does not establish presence or absence. Feature confirmation can coexist with unspecified implementation details.

Derived identifies arithmetic or indexing calculated from the cited definitions. Unmeasured identifies a performance or accuracy claim without a local GPU measurement. Neither label is a substitute for a test result.

A product peak, a source declaration and a runtime result answer different questions. Compiler availability does not establish a library's dispatch or supported deployment. This page does not claim binary compatibility between targets.

Capacity is reported in KiB where the documented bank geometry gives a binary size. Per-CU LDS capacity does not itself establish the maximum allocation for one workgroup in a particular runtime.

2. Matrix instructions

Choose the instruction before choosing the tile layout. A generation label does not fix K, input packing, output register count or sparse metadata. The selected families below are an entry point, not an exhaustive opcode inventory.

Selected matrix instructions; shape is M × N × K, one output block
PropertyCDNA 3 · gfx942CDNA 4 · gfx950
FP16 / BF16 → FP32CDNA 3 · gfx94216×16×16, 32×32×8.CDNA 4 · gfx950Retains these shapes; adds 16×16×32 and 32×32×16.

Clang wider-K signatures

INT8 → INT32CDNA 3 · gfx94216×16×32, 32×32×16.CDNA 4 · gfx950Adds 16×16×64 and 32×32×32.
FP8 / BF8 → FP32CDNA 3 · gfx94216×16×32, 32×32×16; FNUZ inputs.CDNA 4 · gfx950These shapes use OCP inputs; the F8F6F4 family adds wider K and optional block scaling.
Scaled FP8 / FP6 / FP4CDNA 3 · gfx942The gfx950 scaled builtins are not enabled for gfx942.CDNA 4 · gfx95016×16×128, 32×32×64; A and B formats selected independently. E8M0 scale inputs are separate from element data.

Clang scaled signatures

XF32 / TF32 pathCDNA 3 · gfx942Native XF32 MFMA forms include 16×16×8 and 32×32×4.CDNA 4 · gfx950No native XF32 MFMA counterpart in the reviewed ISA. AMD describes TF32 support through software emulation; a library support mark is not an XF32 opcode guarantee.

AMD architecture table and TF32 footnote

Structured sparsityCDNA 3 · gfx942SMFMAC uses compressed A plus selection metadata; FP16/BF16, INT8 and FP8/BF8 families.CDNA 4 · gfx950Retains structured-sparse families and adds wider-K variants. The dense formulas below do not describe compressed A or metadata.

Evidence for the shape rows: CDNA 3 dense instruction table, CDNA 4 dense instruction table, CDNA 4 sparse instruction table, and the pinned LLVM MFMA definitions.

Wider K is an instruction choice

For FP16 inputs and FP32 accumulation, the shared 16×16×16 operation needs 4 FP16 values per lane per input. CDNA 4’s 16×16×32 needs 8. Each still produces 4 FP32 outputs per lane, because M and N did not change. The compiler signatures expose exactly these vector lengths. Common builtin · Wider builtin.

Keeping the original opcode does not automatically use the wider instruction. Larger K may reduce instruction count for a full dot product, but also changes loading, registers and tail handling. Whether that improves a kernel is a measurement question.

Block scaling is part of the operand contract

CDNA 4’s mfma_scale_*_f8f6f4 interface takes packed A/B payloads, format selectors and separate scales. FP4 is E2M1; FP6 can be E2M3 or E3M2. AMD scaled-MFMA explanation.

One scale covers 32 elements along K. For A with shape M×K and B with shape K×N, each A row and B column is divided into 32-element segments, each with its own scale. With K=128, that means four scales per A row and four per B column. This scaling block is distinct from an MFMA output block. ISA §7.2–7.2.1.

The scale byte encodes an exponent. E8M0 uses scale = 2^(b−127) for encoded byte values b=0…254; 0xFF encodes NaN. Thus the byte 0x7F (decimal 127) represents a multiplier of 1, not 127. Scale encoding is separate from the FP8/FP6/FP4 element format. ISA scale encoding.

Packing boundary: Clang’s scaled builtins accept an 8×32-bit vector for each A/B argument, even for smaller payload formats. This interface width is not evidence that every format consumes eight hardware payload VGPRs. Use the selected ISA format’s packing and the emitted code. Clang signatures · ISA §7.1.5.

FP64 and product peak limits

A newer generation does not raise every datatype’s peak. AMD lists matrix FP64 at 163.4 TFLOP/s for MI300X and 78.6 TFLOP/s for MI355X. These are theoretical peaks for those specific products. MI300 table · MI350 table.

They are not measurements of the same kernel, and their ratio does not establish an application speedup or slowdown. Compare the required datatype and actual workload before carrying a performance expectation from one generation to the other.

3. Lane and register layouts

The shared FP16 instruction keeps its output mapping across these generations. In the wider-K CDNA 4 operation, each lane owns more K values while its output rows stay the same.

Let L be a lane index in 0…63, q = floor(L/16) its lane group, and t = L mod 16 its position in that group. e indexes FP16 input elements within one lane; g indexes FP32 accumulator elements. A has shape 16×K, B is K×16, and C/D are 16×16.

Wider K grows the input; the output tile stays 16×16.

Wave64 · one dense FP16 → FP32 block · CBSZ=ABID=BLGP=0. Colors identify q, the group of 16 lanes; every group label is also written in text.

CDNA 3 and 4 · K=16

K=16: four complementary input groupsLane groups 0 to 3 each supply 4 consecutive K positions in A row t or B column t. No group duplicates another.A row t / B column tq=0 · lanes 0–15K 0–3q=1 · lanes 16–31K 4–7q=2 · lanes 32–47K 8–11q=3 · lanes 48–63K 12–154 FP16 values / lane / input

CDNA 4 · K=32

K=32: four complementary input groupsLane groups 0 to 3 each supply 8 consecutive K positions in A row t or B column t. No group duplicates another.A row t / B column tq=0 · lanes 0–15K 0–7q=1 · lanes 16–31K 8–15q=2 · lanes 32–47K 16–23q=3 · lanes 48–63K 24–318 FP16 values / lane / input

Both examples · C/D

The same output rows for K=16 and K=32Each lane group owns four adjacent rows of the 16 by 16 output. Lane position t selects the column; g selects a row within the group.D: 16 rows × 16 columnsq=0 · rows 0–3column tq=1 · rows 4–7column tq=2 · rows 8–11column tq=3 · rows 12–15column t4 FP32 values / lane: D[4q+g, t]

Derived from ISA §7.1.4 and cross-checked with AMD's 16×16×16 FP16 example. This is logical element ownership, not a hardware floorplan.

Scroll horizontally to compare both K sizes. The operand label stays visible.

Derived lane coordinates for the specified FP16 → FP32 operations
OperandK = 16 · CDNA 3 and CDNA 4K = 32 · CDNA 4
A item eK = 16 · CDNA 3 and CDNA 4A[t, 4q + e], 0 ≤ e < 4K = 32 · CDNA 4A[t, 8q + e], 0 ≤ e < 8
B item eK = 16 · CDNA 3 and CDNA 4B[4q + e, t], 0 ≤ e < 4K = 32 · CDNA 4B[8q + e, t], 0 ≤ e < 8
Input packingK = 16 · CDNA 3 and CDNA 4register = floor(e/2)
half = e mod 2
2 payload VGPRs per input.		K = 32 · CDNA 4register = floor(e/2)
half = e mod 2
4 payload VGPRs per input.
C / D item gK = 16 · CDNA 3 and CDNA 4C[4q + g, t] / D[4q + g, t]
0 ≤ g < 4		K = 32 · CDNA 4C[4q + g, t] / D[4q + g, t]
0 ≤ g < 4

These formulas are derived from the ISA’s input and output layout rules with one block, M=N=16, and all broadcast controls zero. Low/high halfwords mean bits 15:0 / 31:16. Register indices here are offsets inside an operand, not allocated physical register numbers. CDNA 4 ISA §7.1.4 · AMD’s shared 16×16×16 FP16 example.

Worked example: lane 18

Here q=1 and t=2. With K=16, lane 18 supplies A[2,4…7] and B[4…7,2]. With K=32, it supplies A[2,8…15] and B[8…15,2]. In both cases its four output values are D[4…7,2]; for example, accumulator item g=2 is D[6,2].

Scope: this is a dense, single-block, FP16 → FP32 example with CBSZ=ABID=BLGP=0. It is not a universal MFMA layout. FP64 output grouping, multiple blocks, broadcast controls, sparse inputs and scaled-format packing require their own definitions. Storage order in global memory is a separate choice from register ownership.

How the coordinate examples were checked

CPU enumeration checks all 64 lanes: A and B each cover 16×K coordinates once, and 64×4 FP32 output items cover every coordinate of the 16×16 result once. This checks the derived indexing, not GPU execution or floating-point accuracy.

The AMD Matrix Instruction Calculator snapshot linked in the sources supports CDNA 3, but does not list CDNA 4. It is not used as proof of CDNA 4 behavior.

4. Numeric formats and cross-lane operations

Audit stored values before tuning instructions. FP8 encoding is a correctness-sensitive change even when an instruction name and buffer size look familiar. Cross-lane operations are a separate question about where a value moves.

FP8 formats

FP8 encodings: identical storage width does not imply identical values
PropertyCDNA 3 · gfx942CDNA 4 · gfx950
E4M3CDNA 3 · gfx942E4M3FNUZ; exponent bias 8; maximum finite magnitude 240; unsigned zero.CDNA 4 · gfx950E4M3FN (OCP); exponent bias 7; maximum finite magnitude 448; signed zero.
E5M2 / BF8CDNA 3 · gfx942E5M2FNUZ; bias 16; no infinity; unsigned zero.CDNA 4 · gfx950E5M2 (OCP); bias 15; infinity and signed zero supported.
HIP device typesCDNA 3 · gfx942__hip_fp8_e4m3_fnuz
__hip_fp8_e5m2_fnuz		CDNA 4 · gfx950__hip_fp8_e4m3
__hip_fp8_e5m2

The device-format split follows HIP low-precision support and ROCm precision support. Bit-format properties are specified in the ONNX float8 reference and AMD’s format table.

Same byte, different finite value.

E4M3 example only. Sign=0, exponent field=7, fraction=0. Values are derived from the format definitions.

CDNA 3 · E4M3FNUZ

CDNA 3: byte 0x38 means 0.5In E4M3FNUZ, exponent field 7 minus bias 8 gives −1; the zero mantissa gives value 0.5.0x38 = 0 | 0111 | 000exponent 7 − bias 8 = −10.5

CDNA 4 · E4M3FN · OCP

CDNA 4: byte 0x38 means 1.0In E4M3FN · OCP, exponent field 7 minus bias 7 gives 0; the zero mantissa gives value 1.0.0x38 = 0 | 0111 | 000exponent 7 − bias 7 = 01.0

Format definitions: ONNX float8. This does not describe rounding, saturation or an automatic conversion between formats.

The byte example is a normal finite value, so it does not depend on saturation or NaN conversion policy. For a real port, inspect the checkpoint/tensor encoding, dequantization scales, conversion mode and special-value handling together. Changing a type name or reinterpreting bytes is not a numerical conversion.

Cross-lane movement

Both ISAs document DS_BPERMUTE_B32 for an indexed gather: each destination lane selects the source lane whose value it receives. DS_PERMUTE_B32 expresses a scatter instead. Their LDS routing hardware does not mean the program must allocate a shared-memory array for the shuffle. CDNA 3 ISA, DS instruction definitions · CDNA 4 ISA, DS instruction definitions.

For a simple gather example with all 64 lanes active and OFFSET=0, aligned byte index 4×7 selects lane 7’s source value. Multiple destination lanes may select it. An inactive source, inactive destination, out-of-range index or colliding scatter is outside that example; consult the exact instruction semantics before relying on its result.

DPP controls describe particular lane patterns and masks; they are not interchangeable with an arbitrary gather. Likewise, an MFMA broadcast modifier changes the matrix operand interpretation rather than acting as a general-purpose shuffle. Keep the matrix’s CBSZ, ABID and BLGP values attached to any layout claim. Common MFMA parameters.

5. Memory and data movement

Revisit both the tile storage and its loading path. CDNA 4 offers more LDS capacity, a different bank count and additional load forms. These are opportunities to retune a kernel, not a prediction that its existing LDS layout will be faster.

LDS and global-to-LDS differences
PropertyCDNA 3 · gfx942CDNA 4 · gfx950
LDS geometryCDNA 3 · gfx94232 × 512 × 4 bytes = 64 KiBCDNA 4 · gfx95064 × 640 × 4 bytes = 160 KiB
Compiler LDS allocation granularityCDNA 3 · gfx942512 bytes at the pinned target definition.CDNA 4 · gfx9501280 bytes at the pinned target definition. Allocations round to this unit; a capacity increase does not imply proportional occupancy.
Global-to-LDS transfer sizesCDNA 3 · gfx942The reviewed direct-load interface supports 1, 2 and 4 bytes per participating lane.CDNA 4 · gfx950Adds 12 and 16 bytes per participating lane. The 12-byte transfer uses a 16-byte destination stride.
LDS read with transposeCDNA 3 · gfx942The gfx950 ds_read_tr* builtin family is not enabled for gfx942.CDNA 4 · gfx950Adds 4-, 6-, 8- and 16-bit transpose variants. These move LDS data into VGPRs; they are not global-to-LDS copies.

Evidence: ISA §2.2.1 for CDNA 3 and CDNA 4; LLVM target features; global-to-LDS reference; Clang transpose-load declarations.

Logical data paths

Direct-to-LDS bypasses one payload staging step.

Logical paths for gfx942 / gfx950. Arrows indicate data movement, not implicit completion or cache bypass.

Staging through VGPRs · both generations

Global memory → VGPR payload → LDS tile → VGPR operands → MFMABoth generations can stage global payload in VGPRs, write LDS, then read matrix operands to registers. Completion and synchronization requirements are omitted from the arrows.Global memoryglobal loadVGPR payloadLDS writeLDS tileLDS readVGPR operandsmatrix operationMFMA

Useful when the kernel must transform loaded values before storing them to LDS.

Direct path · both, with CDNA 4 additions

Global memory → LDS tile → VGPR operands → MFMABoth generations offer direct-to-LDS loading without intermediate VGPR payload staging. CDNA 4 adds 12/16-byte transfers and transpose LDS reads. Later reads still produce register operands.Global memorydirect-to-LDS loadLDS tileLDS read / CDNA 4 transpose readVGPR operandsmatrix operationMFMA

CDNA 4 additions: 12/16-byte direct transfers and element-width-specific LDS transpose reads.

Direct-load completion is tracked as VMEM. A later LDS read and any cross-wave sharing still require the appropriate waits and synchronization.

Conditions and evidence

Direct-load size, destination lane stride and wave-uniform base must match the selected interface. Transpose reads require their own LDS layout; they do not replace arbitrary shuffles or change where the global copy writes.

AMD direct-to-LDS reference · Wider-load instruction gate · Transpose instruction gates · Completion and synchronization.

Direct loads remove the payload’s intermediate VGPR staging. Address registers and other kernel registers still exist. The diagram deliberately omits cache levels and physical wiring. Global-to-LDS semantics.

Copy size is not destination stride

With the immediate offset set to zero, the reviewed interface places lane L at LDS base + L×stride: stride is 4 bytes for 1/2/4-byte transfers and 16 bytes for 12/16-byte transfers. Thus a 12-byte transfer writes three DWORDs in a four-DWORD slot. Packing the next lane at base+12 would describe a different layout. Respect the interface’s wave-uniform LDS base, alignment, bounds and target-specific size constraints. Builtin parameters and destination addressing.

For bank reasoning, a useful address-level calculation is bank = floor(byte_address/4) mod bank_count. Addresses 0 and 128 therefore name the same bank under 32 banks and different banks under 64 banks. This arithmetic alone does not prove a conflict or a speedup: the instruction’s access width, simultaneously serviced lanes and broadcast behavior still matter.

Transpose during an LDS read

The CDNA 4 ds_read_tr* builtin family helps transform a suitable LDS tile into matrix operand fragments. Select the element width and pointer pattern together; one load is not a generic transpose of an arbitrary matrix. In particular, __builtin_amdgcn_ds_read_tr16_b64_v4f16 returns four FP16 values and has a __fp16-based pointer signature at the pinned revision. Clang declaration · AMD transpose-load guide.

Keep this transformation separate from both ordinary ds_read and the wider global-to-LDS copy. A pipeline can use one, both or neither, depending on the required layout.

6. Execution and synchronization

Track the producer and the consumer. Both target families use S_WAITCNT counter-based completion. A barrier coordinates waves in a workgroup; MFMA dependencies require instruction-specific scheduling. They are not substitutes for each other.

Ask what the next operation depends on.

Both targets. These are dependency categories, not three interchangeable instruction sequences.

A consumer needs a value or a shared tile

  1. 1. Is its memory producer complete?

    Global → VGPR or global → LDS: track VMcnt.

    LDS → register: track the relevant LGKMcnt dependency.

    A wait follows the producer's counter, even when a global load writes LDS.

  2. 2. Must other waves rendezvous?

    For a tile shared across waves, complete producer work and synchronize the participating workgroup.

    An ISA S_BARRIER alone does not drain pending memory.

  3. 3. Is it an MFMA register dependency?

    Identify the exact producer and consumer opcodes and the input or accumulator overlap.

    Use MFMA-specific scheduling rules; a memory wait is not a generic arithmetic completion counter.

Independent work may overlap outstanding operations. The required waits depend on the actual instruction stream; this diagram does not prescribe an unconditional full wait after every operation.

LLVM LDS DMA tracking · ISA §4.4 and S_BARRIER · MFMA hazard handling.

Direct-to-LDS needs a memory wait

A direct VMEM-to-LDS operation produces an LDS value without producing a payload VGPR. Waiting only for ordinary LDS operations is therefore insufficient. The LLVM companion explains how the reviewed wait pass tracks pending VMEM-to-LDS writes and later LDS accesses.

For a tile produced and consumed by multiple waves, account for both copy completion and a convergent workgroup synchronization point. At the ISA level, S_BARRIER does not implicitly drain outstanding memory counters. A language-level synchronization primitive may lower to a sequence that includes waits; inspect the emitted code rather than equating the primitive to one machine instruction. CDNA 3 S_BARRIER · CDNA 4 S_BARRIER.

MFMA dependencies need their own analysis

Identify the physical registers and the consumer’s operand role. Here D is a preceding MFMA’s destination; A/B and C are a following instruction’s multiplicands and accumulator. Register overlap is distinct from equal matrix values. MFMA can expose partially written results and can continue reading its inputs after issue. CDNA 3 matrix-core introduction · CDNA 4 matrix-core introduction.

MFMA register dependencies · both generations, with instruction-specific conditions
RelationshipWhat must be checkedReading path
D → A/BRead after writeWhat must be checkedThe next multiplicand overlaps the producer's result. Its wait depends on the producer family and instruction variant.Reading pathSelected LLVM thresholds; accumulator reuse does not cover this relationship.
D → CAccumulator reuseWhat must be checkedDistinguish an exact accumulator chain from partial overlap or a changed matrix instruction.Reading pathFull versus partial C reuse, including short-operation and FP64 exceptions.
MFMA reads C → VALU overwrites CWrite after readWhat must be checkedThe old accumulator input must remain available until the matrix operation has read it.Reading pathInput overwrite and result access; this is separate from memory completion.

The ISA tables distinguish producer families, pass counts and register ranges. Both list a required-wait entry of 2 for a non-dot VALU register write followed by an MFMA read, and 4 for V_CMPX updating EXEC before MFMA. These entries are not a universal wait after every matrix operation. CDNA 3 §7.5, pp.63–66 · CDNA 4 §7.6, pp.74–77.

The LLVM companion cases show selected target-specific thresholds and inspected regression expectations. Keep those compiler rules separate from the ISA contract and from measured instruction latency.

Operand storage matters: these examples do not collapse VGPRs and accumulator registers into one universal rule. Match the actual opcode, register file, register range and modifiers; sparse metadata, FP64 and scaled operations require their own applicable paths. A memory counter does not track MFMA arithmetic completion. For inline assembly, make dependencies visible and audit the instruction sequence.

7. Compiler support and behavior

Compile for the exact processor and verify the emitted path. The full analysis is now in LLVM on CDNA: MFMA dependencies and waits.

Target support and tools

Scope: CDNA 3 gfx942 and CDNA 4 gfx950. The common FP16 builtin is enabled for both; the wider FP16, scaled MFMA and transpose-load interfaces require the reviewed gfx950 features. Default target support and builtin contracts.

Compiler takeaways

  • MFMA spacing depends on the operand role and physical overlap. Exact C reuse, partial C overlap and result reuse as A/B take different paths. Shared FP16 example and conditions.
  • Memory completion is separate. Direct-to-LDS writes need pending-write tracking before later LDS access. Wait-counter handling.
  • Hidden inline-assembly dependencies need an explicit audit. Recognized machine instructions and visible operands define the automatic handling boundary. Developer checks.

Applicability and validation limits

The analysis is pinned to LLVM 4e9d349ab86d. Its wait-state thresholds are not measured GPU latency. Source and regression expectations were inspected; kernel compilation, LLVM regression execution and GPU execution were not performed. Full revision, coverage and limits.

Detailed compiler topics

8. Evidence status and limits

Read instruction inventory, LLVM target support, mechanism descriptions and validation as separate evidence dimensions. The shared status definitions apply throughout this reference; the table below records the evidence found and its limits.

Evidence found and its limits
TopicEvidence foundEvidence limits
Native instruction supportEvidence foundISA: the specified instruction forms. LLVM: the pinned declarations and default target features.Evidence limitsCompatibility with every compiler release, library dispatch or runtime environment.
Register layoutEvidence foundDerived: dense FP16 single-block formulas. CPU checked: coordinate coverage for the illustrated layouts.Evidence limitsAll MFMA shapes, sparse metadata layouts, or a GPU-executed kernel.
MFMA dependenciesEvidence foundSource inspected: ISA dependency categories, pinned LLVM hazard rules and gfx942/gfx950 regression expectations. Selected cases.Evidence limitsLLVM regression execution: Not executed. No compiled-kernel or GPU confirmation; arbitrary inline assembly is outside the automatic-handling claim.
LDS capacity and banksEvidence foundISA: documented geometry. LLVM: allocation granularity.Evidence limitsA universally optimal swizzle or identical service groups for every DS instruction.
PerformanceEvidence foundArchitectural capabilities and clearly labeled product peak examples.Evidence limitsGPU benchmarks: Not executed. Application speedup, measured latency, energy efficiency and optimal tile size remain unmeasured.
NumericsEvidence foundFP8 encodings differ; explicit format selection is required.Evidence limitsGPU accuracy tests: Not executed. Bitwise equality and model accuracy after requantization are not established.

The evidence for these boundaries is linked beside each technical claim in §2–§7. Unspecified mechanisms remain open questions rather than inferred hardware diagrams.

Validation status

Source inspected: the cited ISA, pinned compiler definitions and MFMA hazard implementation, including gfx942/gfx950 regression expectations. CPU checked: coverage of the illustrated dense FP16 coordinate formulas. Not executed: CDNA kernel compilation, LLVM regression tests, GPU execution and benchmarks. Appendix A records these boundaries and the source audits.

Migration priorities

Resolve correctness-sensitive changes before comparing speed. The following order is an engineering recommendation derived from the preceding sections, not an AMD compatibility guarantee.

  1. Identify formats and targets. Separate FNUZ and OCP tensors and confirm the actual processor and compiler.
  2. Select the exact matrix operation. Update input packing, K loops and tails if adopting a wider or scaled instruction; verify sparse metadata separately. Recheck MFMA dependencies when changing opcodes, reusing results as A/B or overlapping accumulator registers.
  3. Audit the memory protocol. Confirm direct-copy destination stride, LDS tile bounds, completion waits and cross-wave synchronization.
  4. Revisit resource use. Check allocated LDS granularity and registers before increasing tile or pipeline depth.
  5. Measure alternatives. Compare the correct baseline with candidate tile, swizzle, transpose-load and copy paths on real hardware.

What a kernel port still needs

  1. Build for gfx942 and gfx950 separately with the compiler version actually used by the project.
  2. Inspect the generated MFMA opcode, operand widths, wait sequence and kernel resource metadata.
  3. Compare outputs against a reference, including tails, strides, special values and quantization scales relevant to the kernel.
  4. Measure the intended shapes on each GPU after correctness passes. Record occupancy, LDS traffic and bank-conflict evidence when evaluating layout changes.

These are the next validation steps; they were not executed as GPU tests for this document. The current checks are listed in the evidence appendix.

9. Sources

The links within each section identify the relevant instruction, table or source lines. These are the main reference entry points.

  1. AMD CDNA 3 ISA, 5 August 2025 — matrix operations, LDS and synchronization semantics.
  2. AMD CDNA 4 ISA, 5 August 2025 — wider/scaled MFMA, register layouts and memory instruction additions.
  3. MI300 microarchitecture and MI350 microarchitecture — product context and theoretical peaks.
  4. AMD Matrix Core Programming, 30 September 2025 — concrete MFMA layout examples and low-precision programming.
  5. ROCm precision support, HIP low-precision types and ONNX float8 specification — format names, representations and interface distinctions.
  6. Global-to-LDS builtins and CDNA 4 transpose-load builtins — programming interfaces; cross-checked against pinned compiler declarations.
  7. LLVM/Clang at 4e9d349ab86d — target bindings, builtin types, instruction gates and wait tracking.
  8. AMD Matrix Instruction Calculator at 94b1f16bf3a2 — a useful CDNA 3 reference; this snapshot does not list CDNA 4 support.

Appendix A. Evidence and validation

Completed: primary-source review; pinned compiler declaration, target-feature and MFMA hazard checks; inspection of gfx942/gfx950 regression expectations; CPU enumeration of the two illustrated dense FP16 layouts; site rendering and browser layout checks. The accompanying source audit records locators and source conflicts; the MFMA compiler audit records dependency formulas, pipeline coverage and test-case boundaries.

Not performed: compilation of a CDNA kernel, LLVM regression execution, GPU execution, or performance/accuracy benchmarking. No claim in this page depends on those tests having passed.

How conflicting or overly broad source descriptions were handled

The general accumulator discussion in the current dense-MFMA handbook can be read as assigning multiple blocks to every packed-input operation. The ISA lists the 16×16×16 FP16 instruction as one block, and the concrete builtin has four FP32 outputs per lane. This page uses those instruction-specific definitions.

The CDNA 4 LDS chapter contains legacy-looking access-count wording alongside its 64-bank description. This page uses the explicit capacity/bank geometry, corroborated by the resource definitions, and does not infer a universal service group or cycle count from that paragraph.

A high-level TF32 support mark can include emulation. The page distinguishes that status from the presence of a native XF32 instruction.

Source review and CPU indexing checks are separate from GPU validation. Diagrams illustrate logical ownership and data paths; they do not assert undocumented physical wiring. See the RDNA comparison for its separate target and wave-size conditions.