Skip to main content

RDNA 3, RDNA 3.5, and RDNA 4

Kernel and compiler reference

A factual comparison of architecture features relevant to GPU kernels and compilers.

Architecture sources rechecked: November 12, 2026; LLVM WMMA handling: November 13, 2026. This is an independently organized reference based on official documentation, not an AMD publication.

Quick comparison

Covered targets: RDNA 3 gfx1100–gfx1103, RDNA 3.5 gfx1150–gfx1153, and RDNA 4 gfx1200/1201. The links below lead to the full conditions and sources.

  • WMMA input fragments change on RDNA 4.

    For wave32, dense 16×16×16 WMMA with FP16/BF16 A/B, each input operand uses 16 elements / 8 VGPRs per lane on RDNA 3 / 3.5, versus 8 elements / 4 VGPRs on RDNA 4. Replicated lane-group inputs become complementary K positions. Operand comparison.

  • The same output count has a different layout.

    With FP32 C/D in that wave32 configuration, all three generations have 8 outputs per lane. RDNA 3 / 3.5 lane groups own alternating rows; RDNA 4 groups own rows 0–7 and 8–15. Formulas and lane examples.

  • Memory completion and WMMA spacing are separate rules.

    Global / buffer loads use S_WAITCNT / VMcnt on RDNA 3 / 3.5 and S_WAIT_LOADCNT / LOADcnt on RDNA 4. The counter follows the memory producer; it does not track WMMA arithmetic completion. Counter families · Register-overlap rules.

  • Feature confirmation and implementation details are separate.

    AMD confirms RDNA 3.5 LPDDR5 awareness, improved compression, and single-use VGPR write suppression. The cited slide leaves implementation details unspecified; cross-generation equivalence is not established by those feature names. Support and evidence limits.

1. Scope and source revisions

This comparison follows the specific targets and source revisions below. Use the target list when applying a result to a kernel: a generation label alone does not establish coverage of every processor in that family.

Default WMMA example: fragment counts and coordinate examples use wave32, 16×16×16 dense WMMA, FP16/BF16 inputs, and FP32 accumulation. The 16-bit packing subsection and sparse-operation rows state their different configurations locally.

AMD’s RDNA documentation index links separate ISA references for RDNA 3, RDNA 3.5, and RDNA 4. This page organizes facts from those manuals and the official programming references. AMD documentation index.

ItemRDNA 3RDNA 3.5RDNA 4
Targets covered by this pagegfx1100–gfx1103.

AMD target list

gfx1150–gfx1153.

AMD target list

gfx1200, gfx1201.

AMD target list

ISA revision checked15 August 2023.

ISA p.1

23 July 2024.

ISA p.1

7 April 2025.

ISA p.1

The target list defines this page’s scope; it is not an exhaustive family inventory. Product-specific cache capacity, memory bandwidth, clocks, platform/driver support, and GFX12.5 are outside this comparison.

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.

Comparison cells: a merged cell shares only the stated property or evidence result, not complete ISA or microarchitectural equivalence. All applicable generations are named, and generation-specific references remain labeled.

Citation placement: cell-level links support that comparison. Generation-heading links support entries in their column, including merged cells, NaN results, layout formulas, and derived examples. Evidence-column links in Section 8 support the entire row.

Source roles: AMD sources describe hardware and programming interfaces; LLVM sources describe compiler targets and semantics. Builtin availability, ISA behavior, and presentation-level claims remain separate evidence categories.

Native-instruction availability uses the default features of the exact targets at the pinned revision. Forced feature overrides, software implementations, and GPU execution tests are outside that check. The three such comparisons also report the AMD XML result; see the evidence snapshot.

These distinctions are collected in the support-status and unresolved-details table.

PDF page numbers throughout this page are physical file positions, counted from 1, rather than printed page numbers.

2. Matrix instructions

For the default wave32 dense WMMA configuration, RDNA 4 halves the per-lane A/B fragment size while retaining eight FP32 C/D values per lane. Check the operand interface first, then register overlap and memory readiness; the sections below keep these separate requirements explicit.

Operands and instruction interfaces

The first three rows describe per-lane storage for wave32, dense 16×16×16 WMMA with FP16/BF16 A/B; each C/D row identifies its accumulator type. For 16-bit C/D, A/B and C/D use the same format: all FP16 or all BF16.
FeatureRDNA 3RDNA 3.5RDNA 4
A/B fragment per lane
RDNA 3 and RDNA 3.5

16 elements · 8 VGPRs per operand.

Replication: lanes t and t+16 supply identical fragments (t=0…15).

RDNA 3: ISA p.84 RDNA 3.5: AMD: A/B layout

8 elements · 4 VGPRs per operand.

Partition: the two lane groups hold complementary K positions.

ISA pp.101–102

FP32 C/D per lane
RDNA 3 and RDNA 3.5

8 outputs.

Lane groups own alternating even and odd rows.

RDNA 3: ISA p.84 RDNA 3.5: AMD: C/D layout

8 outputs.

Lane groups own rows 0–7 and 8–15.

ISA p.102

FP16/BF16 C/D packing
wave32, dense 16×16×16
RDNA 3 and RDNA 3.5

8 results · 8 VGPRs. One selected halfword per register.

OPSEL[2]: 0 selects low; 1 selects high, for both C reads and D writes.

RDNA 3: ISA p.75 · FP16 builtin RDNA 3.5: ISA p.77 · FP16 builtin · BF16 builtin

8 results · 4 VGPRs. Two results per register; even-indexed results use the low halfword.

The corresponding dense builtin has no opsel argument.

ISA p.104 · FP16 builtin

Dense input formats and shapes
RDNA 3 and RDNA 3.5

Listed dense shapes: 16×16×16, with FP16, BF16, IU8 or IU4 A/B inputs.

Listed combinations: FP16 or BF16 inputs use matching A/B formats and either FP32 C/D or the same 16-bit C/D format. IU8/IU4 inputs use I32 C/D.

RDNA 3: ISA p.82 RDNA 3.5: ISA p.84

Also includes FP8 (E4M3), BF8 (E5M2), mixed FP8/BF8 inputs, and 16×16×32 IU4.

ISA p.100 · FP8/BF8 formats

Sparse matrix instruction availability
RDNA 3 and RDNA 3.5

Targets: RDNA 3 gfx1100–gfx1103; RDNA 3.5 gfx1150–gfx1153.

LLVM: Not enabled as native instructions with default target features at revision 4e9d349ab86d.

AMD XML: Not listed. No corresponding instruction entries in the reviewed RDNA 3 / 3.5 files.

XML absence does not establish hardware absence. Evidence scope.

Source links

RDNA 3: AMD XML · LLVM: targets · LLVM: SWMMAC RDNA 3.5: AMD XML · LLVM: targets · LLVM: SWMMAC

Targets: gfx1200/1201.

LLVM: Enabled as native instructions with default target features at revision 4e9d349ab86d.

AMD XML: Listed. Corresponding instruction entries appear in the reviewed file.

ISA p.100 · AMD XML · LLVM: targets · LLVM: SWMMAC

Wave32 builtin names
RDNA 3 and RDNA 3.5
The shared AMD RDNA 3 / 3.5 guide uses the same wave32 builtin family, with names ending in _w32.

RDNA 3: Naming convention RDNA 3.5: AMD target list · Naming convention

The corresponding names end in _w32_gfx12.

Naming convention

RDNA 4 sparse matrix semantics: SWMMAC uses 2:4 sparse A, dense B, and position metadata. For FP16 16×16×32, logical K=32 and stored A has K=16. ISA p.100 · ISA p.106.

These naming examples describe wave32 interfaces. Clang also documents GFX11 wave64 and GFX12 wave64 WMMA builtins.

Dependencies and waits: a shared map for §§2, 6, and 7

Read each applicable branch. A single instruction sequence may need more than one of these checks. Connectors classify requirements; they do not show elapsed cycles, instruction order, or LLVM pass order.

Check the producer and synchronization scope

Keep the requirements separate: memory counters do not track WMMA arithmetic completion, and V_NOP does not complete a memory read. When dependencies coexist, satisfy each applicable rule.

S_DELAY_ALU is optional; mandatory spacing and ISA stall conditions still apply. LLVM behavior is bounded by the reviewed revision and instruction coverage, including inline assembly and expert scheduling.

Register dependencies

Register overlap determines the spacing rule

Scope: the ISA register-dependency rules for the covered RDNA 3 / 3.5 WMMA instructions and RDNA 4 WMMA / SWMMAC instructions. D belongs to the preceding instruction; A/B/C belong to the following instruction. RDNA 4 also checks the following sparse-index input.

Blue shading marks shared VGPR locations. Bars illustrate range relationships; their lengths do not prescribe operand sizes. Arrows in the instruction sequence show program order.

D → A/B overlap
D overlaps a following A/B source The preceding D and following A/B ranges partially overlap. The common register locations are highlighted in blue. Bar lengths are schematic; no operand size or register allocation is prescribed.
  1. Producer matrix instructionWrites the D register range.
  2. Required separatorV_NOP or an independent VALU instruction.
  3. Following matrix instructionReads an A/B range overlapping that D.

Identical or partially overlapping ranges trigger the rule. RDNA 4 also applies it when D overlaps the following sparse-index range.

RDNA 3 / 3.5: one intervening instruction.
RDNA 4: at least one intervening instruction.

D → C only
D overlaps only C The preceding D and following C use the same illustrated register range; other sources are disjoint. This example satisfies the C-only condition. Bar lengths are schematic; no operand size or register allocation is prescribed.
  1. Producer matrix instructionWrites the D register range.
  2. Following matrix instructionReads an overlapping C range; A/B and any RDNA 4 sparse-index input are disjoint from D.

No additional V_NOP gap is specified for this relation. The illustration uses identical D/C ranges; the C-only condition also covers partial D/C overlap.

Hardware stalls may still occur. The table below retains the generation-specific range, instruction-type, and SRC2-modifier conditions. The RDNA 3 source gives no separate D → C stall table or fixed dependency latency.

Combined overlap: if D overlaps both C and A/B (or the RDNA 4 sparse-index input), the required-separator rule still applies. A scalar S_NOP or memory-counter wait is not a specified substitute for that separator. Memory completion is a separate dependency; V_NOP does not establish that a memory read has completed.

Sources

Derived instruction-order schematic. RDNA 3: ISA p.82 · RDNA 3.5: ISA p.86 · RDNA 4: ISA p.100 · ISA p.101.

Dependencies below concern VGPR ranges between matrix instructions. D → C-only excludes overlap of the preceding D with the following A/B or, on RDNA 4, sparse index input.
FeatureRDNA 3RDNA 3.5RDNA 4
D → A/B overlap
RDNA 4: also sparse index
RDNA 3 and RDNA 3.5

Required: one intervening V_NOP or independent VALU instruction.

Trigger: the preceding WMMA D range is identical to or overlaps the following WMMA A/B range.

RDNA 3.5’s scheduling table explicitly classifies this as a correctness requirement. The shared rule specifies instruction separation, not equal WMMA latency or throughput.

RDNA 3: ISA p.82 RDNA 3.5: ISA p.86

Required: at least one intervening V_NOP or independent VALU instruction.

Trigger: the preceding matrix instruction’s D range is identical to or overlaps the following A/B or sparse index range. Applies to WMMA and SWMMAC correctness.

ISA p.100 · ISA p.101

D → C-only: explicit gap
RDNA 3, RDNA 3.5, and RDNA 4

The cited scheduling sections specify no additional V_NOP gap for D → C alone.

Scope: the preceding D does not overlap the following A/B or, on RDNA 4, the sparse index input.

RDNA 3’s stated one-V_NOP condition concerns A/B overlap. RDNA 3.5’s table classifies D → C cases as hardware stall behavior.

The common point is the absence of a stated software gap, not the absence of hardware stalls. Generation-specific stall conditions remain in the next row.

RDNA 3: ISA p.82 RDNA 3.5: ISA p.86 RDNA 4: ISA p.101

D → C-only: hardware stallsThe cited section gives no separate D → C stall table or fixed dependency latency.

ISA p.82

Identical D/C register ranges: stall if WMMA types differ or the consumer uses a SRC2 input modifier (IMOD).

Overlapping D/C register ranges: may stall.

ISA p.86

Identical D/C register ranges: stall if WMMA types differ or the consumer uses SRC2 ABS/NEG.

Overlapping D/C register ranges: may stall.

ISA p.101

The RDNA 3.5 scheduling section also mentions SWMMAC. That wording alone does not establish instruction availability for the targets on this page; the sparse-instruction row separately reports the AMD XML and pinned LLVM target evidence.

WMMA spacing and operand readiness

The rules here describe the ISA contract. The LLVM handling reference separately identifies automatic compiler actions, their conditions, and their limits at a fixed source revision.

Register dependency: D → A/B and D → C refer to the preceding instruction’s destination and the following instruction’s sources, not to equality of matrix values. If D overlaps both C and A/B, the D → A/B correctness rule still applies; RDNA 4 also includes the sparse index input. Absence of an explicit D → C gap does not imply zero latency or stall-free execution. RDNA 3 · RDNA 3.5 · RDNA 4.

V_NOP: a vector no-op, not a memory-completion test. RDNA 4 explicitly defines it as a single-cycle bubble in the vector ALU pipeline. The WMMA overlap rule names V_NOP or an independent VALU instruction as the separator; a scalar S_NOP or scalar memory-counter wait is not one of those specified substitutes. RDNA 3 V_NOP · RDNA 3.5 V_NOP · RDNA 4 V_NOP; the overlap-rule sources are in the table above.

Memory-produced operands: the input read has its own completion dependency. RDNA 3 / 3.5 S_WAITCNT and RDNA 4 S_WAIT_*CNT test outstanding-operation counters against thresholds. Use the producer-specific counter in the completion-wait table; the shared map explains why a WMMA separator cannot replace that wait. RDNA 3 §§5.6–5.7, pp.52–53 · RDNA 3.5 §§5.6–5.7, pp.53–54 · RDNA 4 §5.7, pp.61–65.

S_DELAY_ALU: the manuals separately describe this as optional ALU dependency scheduling that can avoid pipeline stalls; it is not required for correct operation. This statement does not remove the explicit WMMA overlap rule above. RDNA 3 §5.7 · RDNA 3.5 §5.7 · RDNA 4 §5.8.

3. WMMA lane and register layouts

For wave32 dense 16×16×16 WMMA with FP16/BF16 A/B and FP32 C/D, RDNA 4 changes both the input split and the output row mapping. The diagrams show the overall arrangement; the formulas and worked example locate a particular value within it.

The equations transcribe the cited AMD ISA diagrams and fragment-layout definitions. They describe logical matrix coordinates, not memory addresses or instruction timing.

WMMA layout at a glance

Scope: wave32, dense 16×16×16 WMMA, FP16/BF16 A/B inputs, FP32 C/D. Follow one lane pair: t and t+16, with t=0…15.

q=0 · lane t (0–15)q=1 · lane t+16 (16–31)

Inputs show K positions for both A[t,k] and B[k,t]: each filled block holds four values; empty blocks are absent from that lane. Each numbered output tile is a C/D[row,t] row coordinate. All three generations have 8 FP32 outputs per lane.

RDNA 3 / RDNA 3.5

Input K positions · replicated
RDNA 3 / 3.5: replicated input fragments Both lane t and lane t+16 hold all sixteen K positions, 0 to 15. The two lanes supply identical fragments, with sixteen elements in eight VGPRs per input operand. The same K coordinates apply to A[t,k] and B[k,t].

16 elements · 8 VGPRs per lane, per A/B operand.

Both lanes hold the same K=0–15 fragment.

FP32 C/D rows · alternating
RDNA 3 / 3.5: alternating FP32 C/D rows At column t, lane t owns even rows 0, 2, 4, 6, 8, 10, 12, 14, and lane t+16 owns odd rows 1, 3, 5, 7, 9, 11, 13, 15. Each numbered tile is a row coordinate, with eight outputs per lane. C and D use the same mapping.

q=0: even rows 0, 2, …, 14.
q=1: odd rows 1, 3, …, 15.

RDNA 4

Input K positions · partitioned
RDNA 4: partitioned input fragments Lane t holds K positions 0 to 3 and 8 to 11; lane t+16 holds K positions 4 to 7 and 12 to 15. Each lane holds eight elements in four VGPRs per input operand. The same K coordinates apply to A[t,k] and B[k,t].

8 elements · 4 VGPRs per lane, per A/B operand.

The two lanes hold disjoint K positions.

FP32 C/D rows · contiguous halves
RDNA 4: contiguous FP32 C/D rows At column t, lane t owns rows 0 to 7, and lane t+16 owns rows 8 to 15. Each numbered tile is a row coordinate, with eight outputs per lane. C and D use the same mapping.

q=0: rows 0–7.
q=1: rows 8–15.

Tiles follow matrix-coordinate order. The formulas below specify VGPR offsets; for t=0, the diagram matches the lane 0 / lane 16 examples.

Sources

Derived schematic, using the same sources as the formulas: RDNA 3: ISA p.84 · RDNA 3.5: AMD fragment layouts · RDNA 4: ISA p.101 · ISA p.102.

Scroll horizontally to compare the formulas and examples.

WMMA coordinate formulas

All register indices here are offsets from that operand’s starting VGPR.

  • Lane coordinates: L=0…31; q = floor(L / 16) selects the 16-lane group, and t = L % 16 is the position within that group.

  • A/B input storage: element index e advances by VGPR offset, then low/high halfword within that register. h=0 is bits 15:0; h=1 is bits 31:16. The table gives each generation’s range of e.

  • FP32 C/D storage: g=0…7 is the offset of a full 32-bit VGPR holding one accumulator or result.

Each generation column is sourced by the official references in its heading.
MappingRDNA 3 / RDNA 3.5
ISA p.84 · AMD RDNA 3 / 3.5 layouts		RDNA 4
ISA p.101 · ISA p.102
Input element indexe = 0…15e = 0…7
Input storageVGPR offset floor(e / 2), halfword e % 2.VGPR offset floor(e / 2), halfword e % 2.
K coordinatek = ek = 4q + 8 floor(e / 4) + (e % 4)
Input AA[t, k]A[t, k]
Input BB[k, t]B[k, t]
Input relationship between lanes t and t+16Identical A values and identical B values at every input slot e.Disjoint K positions at corresponding input slots; the two lanes together cover K=0…15.
FP32 C / D at VGPR offset gC[2g + q, t]
D[2g + q, t]		C[8q + g, t]
D[8q + g, t]

Worked example: lane 16, FP32 output at VGPR offset 1

Use the wave32 dense 16×16×16 FP16/BF16-input, FP32-output configuration above. Fix L=16 and g=1:

  1. Locate the lane: q = floor(16 / 16) = 1 and t = 16 % 16 = 0.

  2. RDNA 3 / 3.5: substitute into D[2g + q, t]. The row is 2×1 + 1 = 3, so this register holds D[3, 0].

  3. RDNA 4: substitute into D[8q + g, t]. The row is 8×1 + 1 = 9, so the same register offset holds D[9, 0].

Lane and register offset are unchanged; the logical output row changes. C uses the same mapping. This is a substitution into the sourced coordinate formulas, not a separate measurement. The next table expands the result to all eight FP32 registers in lanes 0 and 16.

Example: lanes 0 and 16, with FP32 C/D

Scroll horizontally to compare the formulas and examples.

Derived examples: substitution of L=0 and L=16 into the formulas above.
Lane and operandRDNA 3 / RDNA 3.5
ISA p.84 · AMD RDNA 3 / 3.5 layouts		RDNA 4
ISA p.101 · ISA p.102
Lane 0: ARow 0, K=0–15.Row 0, K={0–3, 8–11}.
Lane 16: ARow 0, K=0–15.Row 0, K={4–7, 12–15}.
Lane 0: BColumn 0, K=0–15.Column 0, K={0–3, 8–11}.
Lane 16: BColumn 0, K=0–15.Column 0, K={4–7, 12–15}.
Lane 0: FP32 DColumn 0, rows {0, 2, 4, 6, 8, 10, 12, 14}.Column 0, rows 0–7.
Lane 16: FP32 DColumn 0, rows {1, 3, 5, 7, 9, 11, 13, 15}.Column 0, rows 8–15.

FP16/BF16 accumulator and output packing

This subsection uses wave32, dense 16×16×16 WMMA with FP16 or BF16 C/D. Let u=0…7 be the logical result index within a lane. V[r][h] denotes halfword h of VGPR offset r in the C or D operand. For RDNA 3 / 3.5, s=OPSEL[2]. C uses the same coordinate and storage mapping as D below.

Eight 16-bit results: selected halves versus packed pairs

Scope: wave32, dense 16×16×16 WMMA, with matching A/B and C/D formats: all FP16 or all BF16. The diagram follows one lane; u=0…7 labels its logical results, and V[r] is a VGPR offset within the C or D operand.

Each register is drawn high halfword on the left (bits 31:16), low halfword on the right (bits 15:0). Blue cells hold the selected results. Gray cells mark unselected halfwords, not additional results of this operation.

RDNA 3 / RDNA 3.5

OPSEL[2]=0 · low halfword

RDNA 3 / 3.5: OPSEL[2]=0, low halfword selected Results u=0 through 7 occupy eight registers, one result in the low halfword of each register. Gray halfwords are unselected. High bits 31:16 are drawn on the left; low bits 15:0 on the right. C reads and D writes use this same map.

8 results · 8 VGPRs.
u → V[u][0]

RDNA 3 / RDNA 3.5

OPSEL[2]=1 · high halfword

RDNA 3 / 3.5: OPSEL[2]=1, high halfword selected Results u=0 through 7 occupy eight registers, one result in the high halfword of each register. Gray halfwords are unselected. High bits 31:16 are drawn on the left; low bits 15:0 on the right. C reads and D writes use this same map.

8 results · 8 VGPRs.
u → V[u][1]

RDNA 4

Two results per VGPR

RDNA 4: two 16-bit results per VGPR Results u=0 through 7 occupy four registers: even u in the low halfword, odd u in the high halfword. High bits 31:16 are drawn on the left; low bits 15:0 on the right. C reads and D writes use this same map.

8 results · 4 VGPRs.
u → V[floor(u / 2)][u % 2]

Even-indexed results occupy the low halfword; odd-indexed results occupy the high halfword.

The two RDNA 3 / 3.5 panels are alternative OPSEL settings for the same eight logical results. C reads and D writes use the same selected halfword. Changing OPSEL changes the storage halfword while preserving the result’s matrix coordinates; those coordinates remain in the formula table below.

Sources

Derived storage schematic. RDNA 3: ISA p.75 · RDNA 3.5: ISA p.77 · RDNA 3 / 3.5 register count: FP16 builtin · RDNA 4: ISA p.104 · FP16 builtin.

Scroll horizontally to compare the formulas and examples.

Formulas and substituted examples use the references in each column heading. Register offsets and builtin vector indices are distinct.
Mapping or exampleRDNA 3 / RDNA 3.5
ISA p.75 · ISA p.77 · Lane coordinates · FP16 slots · BF16 slots		RDNA 4
ISA p.102 · ISA p.104 · FP16 vector · BF16 vector
Coordinate and storage of result uD[2u + q, t] in V[u][s]. One result per VGPR; 8 VGPRs per lane.D[8q + u, t] in V[floor(u / 2)][u % 2]. Two results per VGPR; 4 VGPRs per lane.
Active 16-bit builtin vector indices2u+s:
s=0: {0,2,4,6,8,10,12,14}
s=1: {1,3,5,7,9,11,13,15}		u: {0,1,2,3,4,5,6,7}.
Lane 0: first two resultsD[0,0] in V[0][s]; D[2,0] in V[1][s].D[0,0] in V[0][0]; D[1,0] in V[0][1].
Lane 16: first two resultsD[1,0] in V[0][s]; D[3,0] in V[1][s].D[8,0] in V[0][0]; D[9,0] in V[0][1].
Lane 0: final resultD[14,0] in V[7][s].D[7,0] in V[3][1].
Lane 16: final resultD[15,0] in V[7][s].D[15,0] in V[3][1].

For RDNA 3 / 3.5, changing OPSEL selects the storage halfword for the same eight logical results; it does not change their matrix coordinates or select a contiguous half of the 16-element builtin vector. The unselected halfwords are not additional results of that operation.

4. Arithmetic and cross-lane operations

An instruction’s availability and its edge-case behavior answer different questions. Start with the NaN cases and FP8/BF8 target support, then use the cross-lane diagram to compare who selects the source value and how far it can travel; EXEC conditions remain attached to the relevant operations.

FP32 minimum / maximum

One numeric input and one NaN input; CLAMP=0. Rows align results; cells identify the mode or instruction family that produces each result. For the merged RDNA 3 / RDNA 3.5 cells, the references in both generation headings apply to their respective architectures; RDNA 4 references apply only to RDNA 4.
Result / instruction familyRDNA 3

ISA pp.262–263 · ISA p.31

RDNA 3.5

ISA pp.279–280 · ISA p.31

RDNA 4

ISA pp.304–305 · ISA pp.525–526 · ISA p.83
Instruction families
RDNA 3 and RDNA 3.5
V_MIN_F32
V_MAX_F32

Numeric preference:
V_MIN_NUM_F32 / V_MAX_NUM_F32

NaN propagation:
V_MINIMUM_F32 / V_MAXIMUM_F32

Numeric input selected
RDNA 3 and RDNA 3.5

Quiet NaN: selects the numeric input.

Signaling NaN: selects the numeric input when WAVE_MODE.IEEE=0.

Numeric-preference family: selects the numeric input for either a quiet or signaling NaN.

Quiet NaN produced
RDNA 3 and RDNA 3.5

Signaling NaN with WAVE_MODE.IEEE=1: quiets and propagates the NaN.

NaN-propagation family: produces a quiet NaN for a quiet NaN input; quiets and propagates a signaling NaN.

RDNA 4 signaling-NaN exception: with one numeric input, one signaling NaN, and CLAMP=0, either listed instruction family sets TRAPSTS.INVALID. ISA pp.304–305 · ISA pp.525–526.

FP32 → FP8 / BF8

Four packed / stochastic FP32 → FP8/BF8 conversions. LLVM status uses default target features at revision 4e9d349ab86d; XML omissions do not establish hardware absence. Evidence scope · Snapshot details.
CheckRDNA 3RDNA 3.5RDNA 4
Targets
RDNA 3 and RDNA 3.5
gfx1100–gfx1103 (RDNA 3); gfx1150–gfx1153 (RDNA 3.5).		 gfx1200/1201.
LLVM native instructions
RDNA 3 and RDNA 3.5
Not enabled for the four named conversions.		 Enabled for the four named conversions.
AMD XML inventory
RDNA 3 and RDNA 3.5
Not listed. No corresponding instruction entries in the reviewed RDNA 3 / 3.5 files.		 Listed. The four corresponding instruction entries appear in the reviewed file.
Sources
RDNA 3 and RDNA 3.5
RDNA 3 / 3.5 sources

RDNA 3: AMD XML · LLVM: targets · LLVM: gates RDNA 3.5: AMD XML · LLVM: targets · LLVM: gates
RDNA 4 sources

AMD XML · LLVM: targets · LLVM: gates · LLVM: encodings

Instruction names and semantics are listed below.

RDNA 4 instruction semantics for the four conversions above.
PropertyRDNA 4
Round-to-nearest-even

V_CVT_PK_FP8_F32 / V_CVT_PK_BF8_F32

Converts two FP32 values.

ISA pp.528–529

Stochastic rounding

V_CVT_SR_FP8_F32 / V_CVT_SR_BF8_F32

Seed data is supplied by the second input.

ISA pp.528–529

Overflow

FP16_OVFL=1: signed finite saturation.

FP16_OVFL=0: NaN for FP8; signed infinity for BF8.

ISA p.89

Cross-lane operations

Cross-lane movement: who selects, and how far can data travel?

First compare the reachable lane range, then the source of each lane selector.

Examples: all EXEC lanes active; indices in range; ABS / NEG / OMOD = 0. Arrows carry the original 32-bit source value toward its destination.

Source valueDestination result

1. VALU gather: stay in a row, or select from its partner

Wave32 examples use selector 3 for destinations L5 and L21. The routing applies to the RDNA 3 / 3.5 scalar forms and the additional RDNA 4 _VAR forms compared here. For the scalar examples, FI and BOUND_CTRL are both zero.

Within each 16-lane row

V_PERMLANE16_B32
V_PERMLANE16_VAR_B32

Gather within two separate 16-lane rows Source lane 3 sends x3 to destination lane 5 in row 0. Source lane 19 sends x19 to destination lane 21 in row 1. Each destination selects offset 3 within its own row. Source SRC0Destination Row 0 · lanes 0–15 Row 1 · lanes 16–31 L3: x3 L5: x3 L19: x19 L21: x19

src_lane = 16 × row + selector

From the paired 16-lane row

V_PERMLANEX16_B32
V_PERMLANEX16_VAR_B32

Gather from the paired row Lane 5 selects source lane 19 from row 1. Lane 21 selects source lane 3 from row 0. Arrows carry x19 and x3 across the row boundary. Source SRC0Destination Row 0 · lanes 0–15 Row 1 · lanes 16–31 L3: x3 L5: x19 L19: x19 L21: x3

src_lane = 16 × (row XOR 1) + selector

row = floor(destination_lane / 16). In wave64 the same rule also pairs rows 2 ↔ 3; X16 does not select from an arbitrary row.

Selector difference: the RDNA 3 / 3.5 scalar forms share a 16-slot pattern; RDNA 4 VAR forms use a selector from each destination lane’s VGPR.

2. Selector-source examples: shared pattern vs per-lane VGPR

2. Selector source: shared scalar pattern or per-lane VGPR

RDNA 3 / 3.5 · scalar forms

SRC1/SRC2 hold sixteen 4-bit selectors. L5 and L21 use the same slot, s[5], because both have row offset 5.

s[5] = 3→ same-row gather
L3 → L5andL19 → L21
Sources

RDNA 3 pp.432–433 · RDNA 3.5 pp.458–459

RDNA 4 · additional VAR forms

Each destination reads its own SRC1 VGPR’s low four bits. Matching row offsets can choose different source offsets; the row-routing rules above still apply.

SRC1[L5] & 15 = 3→L3 → L5
SRC1[L21] & 15 = 7→L23 → L21

These two examples use the same-row form.

Sources

RDNA 4 pp.513–514

3. DS gather in wave64: the same index reaches different lanes

DS_BPERMUTE_B32 example: destination L40 supplies byte index ADDR = 28 = 4 × 7, with offset0 = 0. Only the selected transfer is drawn.

RDNA 3 / 3.5 · two 32-lane halves

Lane 40 gathers lane 39 within the upper half-wave Index 7 is relative to lane 40's own 32-lane half. It selects source lane 32 plus 7, lane 39. Lane 7 in the other half is not selected. Half 0 · lanes 0–31 Half 1 · lanes 32–63 Not selected L7: x7 L39: x39 L40: x39

32 + 7 = 39: L40 receives x39.

Sources

RDNA 3 pp.137–138 · RDNA 3.5 pp.139–140

RDNA 4 · all 64 lanes

Lane 40 gathers lane 7 across the half-wave boundary Index 7 selects absolute wave lane 7. The x7 value crosses from lanes 0–31 to destination lane 40. Lane 39 is not selected. Wave64 · lanes 0–63 Not selected L7: x7 L39: x39 L40: x7

L40 receives x7 from the other half of the wave.

Sources

RDNA 4 p.163

EXEC and controls: the active-lane examples do not illustrate inactive-source or suppressed-write cases. The original comparison below retains the scalar FI / BOUND_CTRL controls and the RDNA 4 DS_BPERMUTE_FI_B32 exception. See the cited instruction definitions for masked-lane behavior and index wrapping.

Diagram conventions and gather / scatter semantics

Reading the examples: Lk is lane k; xk is its original 32-bit source value. Arrows point from source data to the destination. All lanes are active in EXEC, indices are in range, and ABS / NEG / OMOD are zero. These are independent examples of instruction semantics, not physical wiring or timing diagrams.

Gather versus scatter: the DS panels show DS_BPERMUTE_B32, where the destination selects its source. With DS_PERMUTE_B32, the source selects its destination; the same generation-specific lane ranges apply. These operations use LDS hardware without allocating LDS storage. Wave32 uses 32 lanes on all three generations.

Cross-lane instruction comparison

FeatureRDNA 3RDNA 3.5RDNA 4
VALU 16-lane gather selectors
RDNA 3 and RDNA 3.5

Scalar selection pattern: SRC0 supplies VGPR data; scalar SRC1/SRC2 form a 64-bit pattern containing sixteen 4-bit lane selectors.

V_PERMLANE16_B32 selects within each 16-lane row. V_PERMLANEX16_B32 selects from the paired row: 0 ↔ 1, and also 2 ↔ 3 in wave64.

Controls: OPSEL[0] is Fetch Inactive; OPSEL[1] is BOUND_CTRL. Other OPSEL bits are reserved. The manuals say to set ABS, NEG and OMOD to zero.

RDNA 3: ISA pp.432–433 RDNA 3.5: ISA pp.458–459

Additional per-lane VGPR selectors.

V_PERMLANE16_VAR_B32 and V_PERMLANEX16_VAR_B32 take both sources from VGPRs; each destination lane supplies its own source-lane selector.

ISA p.94 · ISA pp.513–514

LDS permute lane rangewave64
RDNA 3 and RDNA 3.5

Two independent 32-lane halves.

DS_PERMUTE_B32 and DS_BPERMUTE_B32 use LDS hardware without using LDS memory storage; no LDS allocation is required.

RDNA 3: ISA p.137 RDNA 3.5: ISA p.139

All 64 lanes.

DS_PERMUTE_B32 and DS_BPERMUTE_B32 can access the whole wave.

ISA p.163

RDNA 4 additional permute instruction: DS_BPERMUTE_FI_B32 reads source lanes regardless of EXEC; EXEC still controls destination writes. ISA p.163.

5. Memory and data movement

Start with the payload path between global memory, LDS, and VGPRs, then check whether the named instruction or builtin is available for your target. RDNA 4 transpose loads place data in VGPRs; global-to-LDS builtins have their own target conditions. The later tables cover cache controls, addressing, and atomics as separate choices.

Data movement interfaces

Global, LDS, and VGPR: follow the data payload

Follow the payload destination: ordinary staging, transpose loads, and LDS direct loads use different paths.

Targets and required modes are labeled per panel. The dashed route is a builtin category marked unavailable for the covered targets.

Ordinary loads and stores

RDNA 3 / RDNA 3.5 / RDNA 4

Global and LDS loads and stores use VGPR data operands A global load moves data from global memory into VGPRs; a global store moves it back. An indexed LDS store moves VGPR data into LDS, and an indexed LDS load returns LDS data to VGPRs. Following the two downward arrows stages global data into LDS through VGPRs. Global VGPR LDS GloballoadGlobalstore IndexedstoreIndexedload Global and LDS loads and stores use VGPR data operands A global load moves data from global memory into VGPRs; a global store moves it back. An indexed LDS store moves VGPR data into LDS, and an indexed LDS load returns LDS data to VGPRs. Following the two rightward arrows stages global data into LDS through VGPRs. Global VGPR LDS GloballoadGlobalstoreIndexedstoreIndexedload

Staging example: global load → VGPR payload → indexed LDS store. A later indexed LDS load can supply a VGPR consumer.

Staging and synchronization

Each arrow is a separate operation. The diagram does not imply a fused global-to-LDS instruction or eliminate synchronization requirements.

Sources

RDNA 3: Global p.121 · LDS pp.134–135 RDNA 3.5: Global p.123 · LDS pp.136–137 RDNA 4: Global p.146 · LDS pp.160–161

Transpose during a global load

RDNA 4 · gfx1200 / gfx1201

Global transpose loads write their results into VGPRs GLOBAL_LOAD_TR_B128 or GLOBAL_LOAD_TR_B64 reads global-memory data and changes its matrix storage order while loading it into VGPRs. LDS is not the destination of this instruction. Global VGPR Load + transpose Global transpose loads write their results into VGPRs GLOBAL_LOAD_TR_B128 or GLOBAL_LOAD_TR_B64 reads global-memory data and changes its matrix storage order while loading it into VGPRs. LDS is not the destination of this instruction. Global VGPR Load + transpose

GLOBAL_LOAD_TR_B128
GLOBAL_LOAD_TR_B64

EXEC: zero makes the instruction a no-op; any nonzero value must have all bits set, otherwise behavior is undefined. The ISA specifies wave32 / wave64 layouts separately.

The cited GFX12 builtin target set includes gfx1200/1201 and excludes the covered RDNA 3 / 3.5 targets. This interface result does not describe every possible transpose mechanism.

Sources

RDNA 4 ISA pp.152–153 · AMD builtin targets

LDS direct load

RDNA 4 ISA example · CU mode only

DS_DIRECT_LOAD reads LDS and broadcasts into VGPRs The data starts in LDS at the address supplied by M0. DS_DIRECT_LOAD broadcasts the loaded value into VGPRs under its quad-level EXEC write rule. This is an LDS-to-register path, not a global-to-LDS load. LDS VGPR Read + broadcast DS_DIRECT_LOAD reads LDS and broadcasts into VGPRs The data starts in LDS at the address supplied by M0. DS_DIRECT_LOAD broadcasts the loaded value into VGPRs under its quad-level EXEC write rule. This is an LDS-to-register path, not a global-to-LDS load. LDS VGPR Read + broadcast

DS_DIRECT_LOAD reads LDS; M0 supplies its address and data type. The byte address must be DWORD aligned. WGP mode is excluded.

EXEC is applied per quad: if any lane in a quad is enabled, all four receive the result. This panel illustrates the cited RDNA 4 instruction, rather than an availability comparison with earlier generations.

Sources

RDNA 4 ISA §12.4, p.160

Global-to-LDS builtin category

RDNA 3 / 3.5 / 4: “No” for the four named builtins in AMD’s availability table.

Intended builtin route with a separate availability restriction A dashed arrow identifies the global-to-LDS category, which bypasses payload staging in VGPRs. It does not assert that these builtins are available on RDNA 3, 3.5, or 4; AMD marks all four listed interfaces No for those generations. It also does not establish absence of all hardware memory-to-LDS paths. Global LDS Named builtin routeUnavailable via these builtinson the covered targets Intended builtin route with a separate availability restriction A dashed arrow identifies the global-to-LDS category, which bypasses payload staging in VGPRs. It does not assert that these builtins are available on RDNA 3, 3.5, or 4; AMD marks all four listed interfaces No for those generations. It also does not establish absence of all hardware memory-to-LDS paths. Global LDS Named builtin routeUnavailable via these builtinson the covered targets
How to read the dashed route

This category transfers the payload without a VGPR staging step. The dashed route names the interface’s purpose; it is not a supported-path claim for the covered targets.

Evidence boundary: availability of these four named interfaces and their target feature does not determine whether every hardware memory-to-LDS path is absent.

Sources

AMD builtin availability · Original comparison row

Before consuming data: these paths do not imply completion or visibility. Apply the relevant memory-completion and synchronization rules. The DS permutes in §4 use LDS hardware without storing the payload in LDS; they are distinct from the storage transfers drawn here.

Diagram scope and conventions

Scope: the covered RDNA 3, RDNA 3.5, and RDNA 4 targets; each panel states its instruction or interface scope. Boxes identify data locations, and arrows identify transfer direction. Addresses, caches, physical interconnects, latency, and completion waits are not drawn.

Line styles: solid arrows show the described ISA transfers. The dashed arrow names a builtin category whose availability is limited as stated in its panel.

These rows report builtin availability. The global-to-LDS row is limited to the four named interfaces below; it does not establish the absence of every hardware memory-to-LDS path.
FeatureRDNA 3RDNA 3.5RDNA 4
GFX12 global transpose-load builtin availability
RDNA 3 and RDNA 3.5

Documented builtin availability: the covered RDNA 3 / RDNA 3.5 targets are outside the target set of these GFX12 interfaces.

This is a statement about the cited builtin interfaces, not every possible hardware mechanism for transposing data.

AMD: targets

The documented target set includes gfx1200 and gfx1201.

AMD: targets

Global-to-LDS builtin availability
RDNA 3, RDNA 3.5, and RDNA 4

Documented builtin availability: all four named interfaces below are marked “No” for each of these three generations in AMD’s table.

This shared availability result does not establish that every hardware memory-to-LDS path is absent.

AMD: availability

RDNA 4 transpose-load ISA behavior: Global-memory loads transpose data into VGPRs. The ISA defines GLOBAL_LOAD_TR_B128 and GLOBAL_LOAD_TR_B64. ISA pp.152–153.

The availability row reports AMD’s builtin table, not a prohibition on every possible hardware memory-to-LDS path. It refers specifically to __builtin_amdgcn_global_load_lds, __builtin_amdgcn_load_to_lds, __builtin_amdgcn_raw_ptr_buffer_load_lds, and __builtin_amdgcn_struct_ptr_buffer_load_lds. Clang lists vmem-to-lds-load-insts as their required target feature. Global · Generic · Raw buffer · Structured buffer.

Memory-system features

Each column names the documented change and its scope. The cited presentations do not establish a common implementation for the changes grouped in a row.
FeatureRDNA 3RDNA 3.5RDNA 4
LPDDR5-related reportsReported object and scope

DDR5/LPDDR5 support improvement for Radeon 700M in Ryzen 7040.

Confirmed by AMD for that product.

AMD 2023 p.8 · AMD 2024 p.21

RDNA 3.5 LPDDR5 awareness.

Confirmed by AMD as a memory-subsystem improvement.

AMD 2024 p.21

RDNA 4 LPDDR5 awareness.

Support unconfirmed.

AMD 2025 p.3

LPDDR5-related reportsEvidence and limits

Mechanism equivalence unconfirmed: the cited slides do not establish that the Radeon 700M improvement uses the same mechanism as RDNA 3.5’s LPDDR5 awareness.

AMD 2023 p.8 · AMD 2024 p.21

Details unspecified: the cited slide does not define the implementation or operating conditions of RDNA 3.5 LPDDR5 awareness.

AMD 2024 p.21

The reviewed RDNA 4 presentation documents GDDR6 for the RX 9070 XT, without a corresponding LPDDR5-awareness statement.

AMD 2025 p.3

Documented memory compression

DCC (delta color compression) applies to UAV surfaces and shader writes.

AMD dates shader write compression to RDNA 1 and describes improved handling of view formats and TYPELESS resources in RDNA 3.

DCC: slide 12 · Shader writes: slide 13

RDNA 3.5 “improved compression”. Confirmed by AMD as a memory-subsystem change.

Details unspecified: the cited slide does not identify the algorithm or resource types.

AMD 2024 p.21

AMD describes central compression/decompression in the SoC, with the compression algorithms handled entirely in hardware.

AMD 2025 p.18

The compression row concerns memory resources and data paths, separately from the register-write suppression below. The cited presentations do not identify the RDNA 3.5 compression changes as identical to RDNA 3’s DCC behavior or RDNA 4’s centralized scheme. The Radeon 700M memory statement is product-specific.

Cache controls and addressing

FeatureRDNA 3RDNA 3.5RDNA 4
Global-memory cache-control scope
RDNA 3 and RDNA 3.5

Loads: GLC=0 selects CU/workgroup; GLC=1 selects device.

Stores / atomics: device scope in this model.

RDNA 3: ISA pp.45–46 RDNA 3.5: ISA pp.46–47

Field: two-bit SCOPE: CU / SE (shader engine) / DEV / SYS.

CU scope: VMEM threads in one CU, not the whole WGP.

ISA p.50 · ISA p.146

Temporal hints
RDNA 3 and RDNA 3.5

SLC: 0 = regular; 1 = stream (non-temporal), for graphics-client caches. SMEM operations use SLC=0.

DLC: 0 = regular; 1 = non-temporal, for MALL / Infinity Cache where present; ignored if that cache is absent.

These are the documented control meanings; they do not establish identical cache hierarchies or replacement behavior.

RDNA 3: ISA p.45 · ISA p.46 RDNA 3.5: ISA p.46 · ISA p.47

Field: three-bit TH, with load/store-specific policies.

Policies include: regular, non-temporal, and high-priority temporal. Atomic RMW uses TH[1] for the temporal hint.

ISA p.51

Atomic return control
RDNA 3 and RDNA 3.5

Global-memory atomics: GLC=1 returns the pre-operation memory value to a VGPR.

GLC=0 returns nothing.

RDNA 3: ISA p.46 RDNA 3.5: ISA p.47

TH[0]=1: returns the pre-operation value.

TH[0]=0: returns nothing.

ISA p.51

GLOBAL / SCRATCH / FLAT immediate byte offset
RDNA 3 and RDNA 3.5

Field: 13-bit OFFSET.

GLOBAL/SCRATCH: signed.
FLAT: most significant bit forced to zero.

RDNA 3: ISA p.121 RDNA 3.5: ISA p.123

Field: 24-bit signed IOFFSET.

ISA p.146

RDNA 4 deferred atomic scope: TH[2] controls deferred scope for non-returning cascading atomics. Their full scope is realized at a subsequent release of matching or higher scope. ISA p.51.

Packed global atomics

Availability of GLOBAL_ATOMIC_PK_ADD_F16 and GLOBAL_ATOMIC_PK_ADD_BF16, including returning and non-returning forms.
FeatureRDNA 3RDNA 3.5RDNA 4
Packed 16-bit global atomic add
RDNA 3 and RDNA 3.5

Targets: RDNA 3 gfx1100–gfx1103; RDNA 3.5 gfx1150–gfx1153.

LLVM: Not enabled as native instructions with default target features at revision 4e9d349ab86d.

AMD XML: Not listed. No corresponding instruction entries in the reviewed RDNA 3 / 3.5 files.

XML absence does not establish hardware absence. Evidence scope.

Source links

RDNA 3: AMD XML · LLVM: targets RDNA 3.5: AMD XML · LLVM: targets

Targets: gfx1200/1201.

LLVM: Enabled as native instructions with default target features at revision 4e9d349ab86d.

AMD XML: Listed. Corresponding instruction entries appear in the reviewed file.

ISA p.707 · AMD XML · LLVM: targets

LLVM instruction definitions for the availability row: LLVM: F16 · LLVM: BF16 · LLVM: encodings. These pinned predicates and encodings are shared references for the three processor columns above.

RDNA 4 packed atomic-add semantics

AspectRDNA 4 behavior
Packed addition

GLOBAL_ATOMIC_PK_ADD_F16 and GLOBAL_ATOMIC_PK_ADD_BF16 perform component-wise additions of two packed FP16 or BF16 values in a 32-bit memory word.

ISA p.707

Return control

TH[0] controls whether the original value is returned.

TH fields: p.51

Temporal reuse hint

TH[1] separately selects the temporal reuse hint.

TH fields: p.51

Instruction and platform scope

  • Transpose EXEC semantics: EXEC=0 makes the transpose-load instructions no-ops; any other EXEC value must have all bits set, otherwise behavior is undefined. ISA §11.6.2.
  • LDS direct load: DS_DIRECT_LOAD reads LDS into VGPRs and is restricted to CU mode. ISA §12.4.
  • RDNA 3.5 APUs: AMD’s system guide describes per-process GPUVM address spaces for Ryzen APUs with gfx1150, gfx1151, and gfx1152. ROCm 10.1 memory settings.

6. Execution and synchronization

For global/buffer reads, completion uses VMcnt on RDNA 3 / 3.5 and LOADcnt on RDNA 4; the producer instruction class determines the counter. Use this chapter to check register state, memory completion, and workgroup synchronization, with the shared dependency map showing how they relate to WMMA spacing.

Execution and register state

FeatureRDNA 3RDNA 3.5RDNA 4
Wave sizes
RDNA 3, RDNA 3.5, and RDNA 4

Supported wave sizes: wave32 and wave64. A shader program is compiled for and runs at a particular wave size.

Instruction-specific restrictions still apply: VOPD is wave32-only in all three manuals.

RDNA 3: ISA p.18 · VOPD p.77 RDNA 3.5: ISA p.18 · VOPD p.79 RDNA 4: ISA p.19 · VOPD p.93

Scalar floating pointLLVM explicitly marks SALU floating-point instructions unavailable on gfx1100–gfx1103.

LLVM: processors

RDNA 3.5 and RDNA 4

Shared operation categories: FP32/FP16 scalar arithmetic and comparisons, plus conversions.

Operand modifiers: not supported.

FP32 min/max names listed — RDNA 3.5: S_MIN_F32, S_MAX_F32.

RDNA 4: S_MIN_NUM_F32, S_MAX_NUM_F32, S_MINIMUM_F32, S_MAXIMUM_F32. The category summary does not imply identical instruction sets or numerical behavior.

RDNA 3.5: ISA p.62 RDNA 4: ISA p.73

Single-use VGPR writes

Support unconfirmed.

The reviewed RDNA 3 architecture and presentation sources do not explicitly establish suppression of VGPR writes for single-use values.

AMD 2023 p.8 · ISA p.26

Confirmed by AMD: suppression of single-use VGPR writes.

Details unspecified: the cited slide does not describe triggering conditions.

AMD 2024 p.21

Support unconfirmed.

The reviewed RDNA 4 architecture presentation describes dynamic register allocation, without explicitly confirming single-use write suppression. Historical compiler hint support is documented separately below.

AMD 2025 p.10 · Compiler-history evidence

Documented VGPR management
RDNA 3 and RDNA 3.5

Whole-wave release. Through S_SENDMSG, with no later reallocation. After release, the only valid action is to terminate the wave.

RDNA 3: ISA p.26 RDNA 3.5: ISA p.26

Dynamic allocation. Compute waves launched in dynamic-VGPR mode; wave32 only, excluding graphics shaders.

Allocation: S_ALLOC_VGPR; 16- or 32-VGPR blocks. Block size is chip-wide, not per dispatch.

Result: failure SCC=0; success SCC=1.

ISA p.28 · ISA p.29

Documented VOPD pairing semantics
wave32
RDNA 3 and RDNA 3.5
The manuals describe simultaneous execution: a read paired with a write to the same VGPR observes the old value. Other VOPD pairing constraints still apply.

RDNA 3: ISA p.77 RDNA 3.5: ISA p.79

The manual prohibits X from overwriting a Y source. Y’s destination may overlap an X source. Other VOPD pairing constraints still apply.

ISA p.93

Dynamic-VGPR mode in LLVM: the cited documentation dates the mode to GFX12. LLVM: mode.

Completion waits and barriers

The counter rows concern memory completion, not WMMA arithmetic completion. The separate WMMA register-dependency rules distinguish D → A/B from D → C-only; scalar counter waits are not the specified V_NOP / independent-VALU separator. Producer-specific counter details follow this table.
FeatureRDNA 3RDNA 3.5RDNA 4
Wait mechanism
RDNA 3 and RDNA 3.5

S_WAITCNT tests VMcnt, LGKMcnt and EXPcnt against encoded thresholds; each enabled test waits until its counter is at or below the threshold.

No-wait encodings: 0x3f for VMcnt / LGKMcnt, and 0x7 for EXPcnt, disable the corresponding wait.

The counters track outstanding instructions, not threads or a fixed number of elapsed cycles.

RDNA 3: ISA pp.52–53 · ISA p.247 · ISA p.214 RDNA 3.5: ISA pp.53–54 · ISA p.263 · ISA p.221

S_WAIT_*CNT: waits until the selected counter is at or below its encoded threshold.

ISA p.62 · ISA pp.63–65 · ISA p.287

Global / buffer loads
RDNA 3 and RDNA 3.5

S_WAITCNT · VMcnt.

The VMcnt group also includes vector-memory atomics that return data; it is not a counter dedicated only to global / buffer loads.

RDNA 3: ISA pp.52–53 RDNA 3.5: ISA pp.53–54

S_WAIT_LOADCNT · LOADcnt.

Includes regular vector-memory loads and returning atomics.

ISA p.62 · ISA pp.63–65 · ISA p.287

Indexed LDS operations
RDNA 3 and RDNA 3.5

S_WAITCNT · LGKMcnt for indexed LDS operations.

LGKMcnt is also used by SMEM, GDS/GWS, FLAT and messages. Direct / parameter LDS loads instead use EXPcnt.

RDNA 3: ISA pp.52–53 RDNA 3.5: ISA pp.53–54

S_WAIT_DSCNT · DScnt.

ISA p.62 · ISA pp.63–65 · ISA p.287

Vector-memory stores / non-returning atomics
RDNA 3 and RDNA 3.5

S_WAITCNT_VSCNT · VScnt.

In all-in-order mode, VScnt is unused; VMcnt tracks all vector-memory events, including stores and non-returning atomics.

RDNA 3: ISA p.214 · VMcnt: p.215 RDNA 3.5: ISA p.221 · VMcnt: p.222

S_WAIT_STORECNT · STOREcnt.

ISA p.62 · ISA pp.63–65 · ISA p.287

Workgroup barriers
RDNA 3 and RDNA 3.5

Barrier interface: S_BARRIER combines arrival and wait on one implicit workgroup barrier.

LLVM defines memory synchronization separately through paired release/acquire fences, synchronization scopes and address spaces.

LLVM GFX6–11 barriers · LLVM: memory synchronization

Separate S_BARRIER_SIGNAL and S_BARRIER_WAIT; ordinary shader workgroup barrier ID -1.

LLVM GFX12 barriers · LLVM barrier-ID table

Counter exceptions: FLAT instructions use both the vector-memory and LDS-related counters; direct/parameter LDS loads use EXPcnt rather than the indexed-LDS counter. These are memory-producer rules, separate from WMMA register spacing. RDNA 3 counter groups · RDNA 3.5 direct/parameter loads · RDNA 3.5 FLAT · RDNA 4 FLAT and LDS.

Single-use VGPR writes: evidence boundary

AMD statementAMD’s Hot Chips 2024 presentation confirms the RDNA 3.5 feature but does not specify triggering conditions. AMD 2024 p.21.

Historical implementationA 2023 LLVM contribution authored by AMD engineers described S_SINGLEUSE_VDST annotations and conditional write suppression: forwarding a single-use VALU result before VGPR writeback could avoid that write.

Compiler targets and removalLLVM modeled this hint for GFX11.5 and GFX12 (gfx1200/1201 processor bindings); in September 2024, LLVM removed the unused compiler feature definition, instruction definitions, and pass.

Hardware evidence boundaryThese compiler records leave RDNA 3 / RDNA 4 hardware support unconfirmed. They also do not establish the full connection between the historical hint and the feature reported in AMD’s RDNA 3.5 presentation. Removing compiler support does not establish hardware absence.

Instruction and memory-model details

  • Legacy opcode: RDNA 4’s S_WAITCNT machine opcode is equivalent to S_WAIT_IDLE and ignores its operand. ISA p.281.
  • Barrier visibility: LLVM’s memory model pairs barrier arrival with a preceding release fence and barrier wait with a following acquire fence; synchronization depends on scopes and address spaces. Execution-barrier memory model.
  • Later extensions: LLVM lists named barrier IDs 0–16 and cluster barrier IDs -4/-3 as GFX12.5 features, outside this page’s target scope. LLVM barrier-ID table.

7. Compiler support and behavior

Use the exact target and inspect the emitted sequence. The full analysis is now in LLVM on RDNA: WMMA dependencies and waits.

Target support and tools

Scope: gfx1100–gfx1103, gfx1150–gfx1153, and gfx1200/1201. Native availability remains beside the SWMMAC, FP8 conversion, and packed atomic comparisons. Compiler targets and ISA tools gives the supporting interfaces.

Compiler takeaways

Applicability and validation limits

The analysis is pinned to LLVM 4e9d349ab86d and recognized machine instructions in the reviewed pipelines. Source and regression expectations were inspected; LLVM regression tests 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 focuses on feature reports and unresolved mechanisms.

Feature reports and unresolved mechanisms

Three entries have unconfirmed support: RDNA 4 LPDDR5 awareness and RDNA 3 / RDNA 4 single-use VGPR write suppression.

AMD confirms the three RDNA 3.5 features listed here; their remaining gaps concern implementation details.

The Radeon 700M memory improvement is confirmed for that product, while its equivalence to RDNA 3.5 LPDDR5 awareness is unconfirmed. Compression equivalence across generations is also unresolved.

An unspecified implementation is not an unconfirmed feature. These labels describe the reviewed sources; see the status definitions.

Official links in “Evidence found” apply to the whole row. PDF links identify physical file pages; LLVM records describe compiler history separately from hardware status.
Item and scopeEvidence statusEvidence foundEvidence limits
RDNA 3: equivalence to RDNA 3.5 LPDDR5 awareness

Radeon 700M improvement: confirmed by AMD.

Mechanism equivalence: unconfirmed.

AMD reports improved DDR5/LPDDR5 support for Radeon 700M in Ryzen 7040.

Its RDNA 3.5 slide separately names LPDDR5 awareness.

AMD 2023 p.8 · AMD 2024 p.21

The cited slides do not describe a shared mechanism or explicitly map the Radeon 700M improvement to RDNA 3.5’s LPDDR5-awareness feature.

The Radeon 700M confirmation is product-specific.

RDNA 3.5: LPDDR5-awareness implementation

LPDDR5 awareness: confirmed by AMD.

Implementation details: unspecified.

AMD lists LPDDR5 awareness among the memory-subsystem improvements.

AMD 2024 p.21

The cited slide does not define the mechanism, the behavior changed by “awareness,” or its operating conditions.
RDNA 4: LPDDR5 awareness

Support unconfirmed.

The RDNA 4 presentation specifies GDDR6 for RX 9070 XT.

The explicit LPDDR5-awareness statement is in the RDNA 3.5 presentation.

AMD 2025 p.3 · AMD 2024 p.21

No explicit RDNA 4 statement confirming or denying this mechanism was identified in the reviewed architecture material.

RX 9070 XT’s GDDR6 interface does not establish whether RDNA 4 implements LPDDR5 awareness.

RDNA 3.5 compression: details and relationship to RDNA 3 / 4

Compression: confirmed by AMD in all three generations.

RDNA 3.5 details: unspecified.

Mechanism equivalence: unconfirmed.

AMD documents RDNA 3 DCC and shader-write compression, RDNA 3.5 “improved compression,” and RDNA 4 central compression/decompression.

AMD RDNA 3, slide 12 · Slide 13 · AMD 2024 p.21 · AMD 2025 p.18

The RDNA 3.5 slide does not identify algorithms or resource types. The cited decks do not identify its changes as the same scheme as RDNA 3 DCC or RDNA 4 centralized compression.
RDNA 3: suppression of single-use VGPR writes

Support unconfirmed.

The reviewed RDNA 3 register description and Radeon 700M architecture slide do not explicitly describe this suppression.

AMD explicitly reports it for RDNA 3.5.

ISA p.26 · AMD 2023 p.8 · AMD 2024 p.21

No RDNA 3-specific hardware description confirming or denying suppression of these writes was identified in the reviewed sources.

The positive RDNA 3.5 statement does not establish the preceding generation’s behavior.

RDNA 3.5: single-use VGPR suppression conditions

Write suppression: confirmed by AMD.

Triggering conditions: unspecified.

AMD’s slide reports skipping single-use VGPR writes.

The historical AMD-authored LLVM implementation describes forwarding a single-use VALU result before VGPR writeback.

AMD 2024 p.21 · LLVM: historical mechanism · LLVM removal

The presentation does not specify triggering conditions. The reviewed material does not establish the full connection between that reported feature and the removed compiler hint.

The historical compiler description is not a complete current hardware contract.

RDNA 4: suppression of single-use VGPR writes

Support unconfirmed.

AMD’s presentation describes dynamic VGPR allocation.

Historical LLVM support modeled a single-use hint for GFX12; those compiler definitions and the pass were later removed.

AMD 2025 p.10 · LLVM: historical hint · gfx1200/1201 bindings · LLVM removal

No explicit AMD hardware statement resolving RDNA 4 single-use write suppression was identified. The LLVM removal does not supply a hardware-support conclusion.

Documented dynamic allocation and historical compiler support do not establish single-use write suppression in current RDNA 4 hardware.

Validation status

Source inspected: the ISA references and XML inventory identified in Appendix A, plus the pinned compiler code and regression expectations in the LLVM companion evidence section. Not executed: those LLVM regression tests and GPU benchmarks. The linked evidence records the exact source-check coverage; source inspection does not establish GPU behavior or performance.

9. Sources

This list identifies the source documents; precise page and source-line links remain beside the claims. See citation and status conventions for how those links apply, and the evidence appendix for the pinned snapshots and checks.

  1. AMD RDNA architecture documentation

    Official entry point for the generation-specific ISA references.

  2. RDNA 3 Instruction Set Architecture Reference Guide

    Execution state, instruction semantics, registers, memory operations, and dependencies.

  3. RDNA 3.5 Instruction Set Architecture Reference Guide

    The independent RDNA 3.5 contract, including SALU floating-point operations in §6.8.

  4. RDNA 4 Instruction Set Architecture Reference Guide

    Cited sections cover dynamic VGPR allocation, cache controls, instruction dependencies, scalar and vector arithmetic, WMMA, memory addressing, atomics, and cross-lane operations.

  5. AMD ROCm Optimization Guide: RDNA 3 dense WMMA builtins

    Covers RDNA 3 and RDNA 3.5, including target availability, fragments, and builtin interfaces.

  6. AMD ROCm Optimization Guide: RDNA 4 dense WMMA builtins

    Target availability, builtin signatures, and input formats. The RDNA 4 coordinate mappings on this page are linked directly to ISA diagrams.

  7. AMD ROCm Optimization Guide: RDNA 4 WMMA transpose load builtins

    GFX12 transpose builtin target availability. EXEC semantics are sourced to the ISA.

  8. AMD ROCm Optimization Guide: Global-to-LDS builtin availability

    Availability and feature requirements by builtin and target.

  9. AMD Hot Chips 2024 presentation, slide 21

    RDNA 3.5 shader and mobile-memory improvements in AMD's presentation hosted by the conference.

  10. AMD: RDNA 3 — Beyond the Current Gen

    Lou Kramer, Reboot Develop Blue, April 2023. Slides 12–14 describe DCC and shader write compression.

  11. AMD Ryzen 7040 Series Technology Overview — Hot Chips 2023

    Slides 3 and 8 describe the Phoenix memory platform and Radeon 700M memory support.

  12. AMD RDNA 4 Radeon 9000 Series GPU — Hot Chips 2025

    Andy Pomianowski and Laks Pappu, August 2025. Architecture, memory compression, and register-management material.

  13. AMD machine-readable GPU ISA documentation

    Supported architecture list, XML specification downloads, and IsaDecoder tools.

  14. AMD machine-readable ISA archive (2026-08-06 package)

    Checked files, release metadata, entry counts, encoding values, and archive fingerprint are recorded in the XML / LLVM cross-check.

  15. LLVM AMDGPU source, pinned revision

    Processor bindings, feature inheritance, native-instruction predicates, and real encodings. Cell-level links point to the relevant source lines at this fixed revision.

  16. AMD-authored LLVM single-use VGPR hint implementation

    Conditional write-suppression description, November 2023. The September 2024 removal is recorded separately from hardware availability.

  17. LLVM AMDGPU Usage

    LLVM's target, feature, ABI, code-object, and memory-model documentation.

  18. Clang AMDGPU Builtin Reference

    Builtin signatures, operand types, required target features, and explicit wave64 entries.

  19. LLVM AMDGPU Execution Synchronization: GFX12

    Barrier mechanisms and the boundary between GFX12 and later extensions.

  20. AMD RDNA 3.5 system optimization

    Ryzen APU GPUVM and system-memory considerations; platform guidance rather than a WMMA specification.

Appendix A. Evidence snapshots and checks

This appendix records how the comparison was checked. It retains the artifact metadata and source bindings needed to revisit the results alongside the native-availability comparisons.

Compiler revision and regression scope

The LLVM companion evidence section records the full revision and inspected WMMA regression expectations, including their wave-size conditions. These tests were not executed.

Native instruction availability: XML and LLVM cross-check

The SWMMAC, FP32-to-FP8/BF8 conversion, and packed 16-bit global-atomic rows use two separate checks: AMD XML instruction names and encodings, and LLVM’s processor-feature inheritance plus instruction predicates. LLVM status refers to default instruction features for the exact processors in this page’s scope, not forced feature overrides, software implementations of the same operation, or GPU execution tests.

The AMD XML archive was retrieved on November 12, 2026.

The three files amdgpu_isa_rdna3.xml, amdgpu_isa_rdna3_5.xml, and amdgpu_isa_rdna4.xml each declare release date 2026-02-20 and schema version 1.1.1. They contain 1,148, 1,208, and 1,264 instruction entries respectively.

The linked archive is the dated 2026-08-06 package; its SHA-256 is 82404f1126761b7877595b622afa7e1f311f2f41e89a3abe9aaf8ad045c082e2.

LLVM source is pinned to 4e9d349ab86dc0a913db3cba11aa2763d6985aff. RDNA 3 / 3.5 processor bindings and RDNA 4 bindings map the named processors to their feature sets. The shared GFX11 feature set and inherited generation features do not supply the missing instruction gates. The per-target additions are gfx1100–1103 and gfx1150–1153 and gfx1200/1201.

These conclusions do not extend to all GFX11 processors: the same revision enables FP8 conversion and SWMMAC features for gfx1170’s feature set, outside this comparison.

The generated site includes local styles, diagrams, and scripts. External source documents require an internet connection. Document content and the on-page contents remain readable without JavaScript; search and automatic highlighting require JavaScript.

This reference contains no GPU benchmark results and does not claim complete ISA or binary compatibility between generations.