Capability Gain from ISA to Shared Execution on AMD MI300A and MI250X

arXiv:2611.07198v1 [cs.DC] 4 Nov 2026

Capability Gain from ISA to Shared Execution on AMD MI300A and MI250X Aaron Jarmusch

Sunita Chandrasekaran

dept. of Computer Information Sciences University of Delaware Newark, US

dept. of Computer Information Sciences University of Delaware Newark, US schandra@udel.edu

Abstract—Each new GPU generation is showcased by its peak instruction rates, yet users hardly ever run instructions directly: they reach the hardware’s performance through vendor libraries, applications and devices shared with other jobs. How much of a theoretical speedup survives these layers is rarely measured, so a team deciding whether to move to a new generation, or to adopt FP8 or 2:4 sparsity, lacks adequate guidance. AMD’s MI300A (CDNA3) adds new FP8 matrix instructions and 2:4 structuredsparse instructions that the predecessor, the MI250X (CDNA2), does not have. We analyze this theoretical speedup gap on AMD’s MI300A (CDNA3) against the MI250X (CDNA2) by defining a new metric, capability gain. Capability gain is the fraction of a feature’s instruction-level speedup that it keeps at a higher level, such as a library, an application or a shared device. We trace it from ISA microbenchmarks through hipBLASLt and hipSPARSELt to LLM inference and mixed-precision HPL, and onto a device shared by up to eight concurrent streams. FP8 keeps most of its speedup through the library, completing an 81923 GEMM 3.95× faster than the MI250X with FP16, but unfused quantization turns it into a perlayer slowdown in an LLM, and 2:4 sparsity keeps at most about half of its speedup in hipSPARSELt. Under queue sharing, FP8’s GEMM capability gain rises above 1. How fairly streams share the device depends on the number of hardware queues, not on the numeric precision, yet shared LLM serving gains nothing from FP8. Index Terms—AMD MI300A, AMD MI250X, CDNA, FP8, structured sparsity, matrix cores, GPU sharing, microbenchmarking

I. I NTRODUCTION Recent GPU generations increase matrix throughput mostly through narrower formats and structured sparsity. Deep learning uses FP8 [1], [2] and 2:4 sparsity [3], and HPC solvers now use mixed precision [4], [5]. Vendors advertise each feature by its peak instruction rate, but users reach it through libraries and applications. Before adopting FP8, 2:4 or a new generation, users need to know how much of that rate their own code gets. AMD added FP8 and 2:4 structured sparsity to its CDNA3 GPUs. El Capitan, the TOP500’s fastest system from November 2024 to June 2026 [6], [7], is built on CDNA3’s MI300A APU, and Frontier, still in production, is built on CDNA2’s MI250X, which has neither feature. On paper, a CDNA3 compute unit (CU) can do 2× the FP16 matrix work per cycle of a CDNA2 CU, 4× with FP8, and 8× with FP8 and 2:4 sparsity [8]–[10]. AMD does not report how much of these

rates a library, an application or a shared device keeps, or how much faster CDNA3 is than CDNA2 in practice. The closest study, by Hanindhito and Patel [11], characterizes CDNA3 at the instruction, library and application levels on the MI300X and MI325X. They measure FP8 only on the matrix cores, measure neither sparsity nor sharing, and do not compare across generations. The one generational AMD comparison, by Tee et al. [12], covers only memory. As Table II shows, no prior AMD study measures FP8 or 2:4 against the previous generation. We track each capability with one metric, capability gain: the speedup it delivers at a level divided by its instruction-level speedup (Section IV-C). Fig. 1 shows it on the MI300A at four levels: the instruction set (ISA), the vendor library (LIB), the application (APP) and the shared device (SHR), where several GEMM streams run at once. Fig. 1(a) shows that FP8 keeps 0.580–1.01 of its instruction-level speedup in hipBLASLt, but a whole LLM layer in FP8 runs at only 0.862× FP16’s speed. Fig. 1(b) shows that 2:4 keeps at most 0.554 in hipSPARSELt, which is about no speedup. Sharing helps both: with eight streams, FP8’s capability gain rises to 1.23 and sparse runs 1.39× faster than dense. Each drop in Fig. 1 lines up with a specific boundary, such as the instruction clock, the library’s kernel choice, an unfused conversion to FP8 or a shared queue pool, and each has a cause we can measure. We answer four research questions. RQ1: Does the instruction-level speedup survive the library (Sections V and VI)? RQ2: Where does 2:4 lose its speedup (Section VII)? RQ3: How does sharing a device change capability gain (Section VIII)? RQ4: How much of FP8’s and 2:4’s speedup reaches real applications (Section IX)? To our knowledge, this is the first study to measure FP8 and 2:4 on AMD hardware against the previous generation, and the first to measure sharing on both generations under one protocol, with the number of hardware queues (4 or 16) as a controlled variable. In summary, this paper makes the following contributions: • We define capability gain and a method to trace it across four levels (instruction, library, application and shared device) on two generations, under a protocol that controls kernel selection and operand values, and we attribute each loss to a mechanism with hardware counters and measured clocks (Section IV, Fig. 1).

TO

1 3143 16GB

stack

1.0

2

0.5

1

LIB

APP

FP8 over FP16 (a) 2:4, FP16 (b)

SHR

ISA

LIB

2:4, BF16 (b) range

APP

38ComputeUnits

38ComputeUnits

4 MBL2Cache

4 MBL2Cache XCD 38ComputeUnits

4 MBL2Cache

4 MBL2Cache

1 3M

XCD

16GB

38ComputeUnits

8 Zen Cores 32MB 13

4 MBL2Cache

c

1 3M

XCD

8 Zen Cores 32MB 13

16GB

38ComputeUnits

c

4 MBL2Cache

8 Zen Cores 32MB 13

CC

TO

(a)

16GB

stack

1 3M

XCD 38ComputeUnits

Stack

fitted library rate

1 3M

XCD

16GB

Stack

SHR

TO

XCD

1 3M

Stack

ISA

Infinity Fabric

Infinity Cache

(b) 2:4 sparse over dense

speedup (×)

capability gain (speedup / ISA speedup)

(a) FP8 over FP16

16GB

Stack

1 3M 16 Stack

1 3M 16GB

stack

TO

Compute Unit (1 of 228) HIP streams

Fig. 1. Capability gain of FP8 and 2:4 on the MI300A at four levels. Points: GEMM shapes in increasing size (LIB); prefill, decode and one layer (APP); 1, 2, 4 and 8 streams (SHR). Dashed: full speedup; dotted: none.

Hardware queues (4 or 16)

TABLE I A RCHITECTURAL PARAMETERS OF THE MI300A AND MI250X. MI300A (CDNA3) MI250X (CDNA2) Dies CUs Peak engine clock LDS per CU FP16 FLOP/CU/clock FP8 FLOP/CU/clock Peak FP16 matrix Peak FP8 matrix Peak 2:4 FP16 Peak 2:4 FP8

6 XCDs 228 (38 per XCD) 2.1 GHz 64 KiB 2,048 4,096 981 TFLOP/s 1,961 TFLOP/s 1,961 TFLOP/s 3,922 TFLOP/s

2 GCDs 220 (110 per GCD) 1.7 GHz 64 KiB 1,024 – 383 TFLOP/s – – –

We compare CDNA3’s FP8 matrix path on the MI300A with CDNA2’s FP16 path on the MI250X, from instruction microbenchmarks through hipBLASLt, per package and per CU·clock (Sections V and VI). • We evaluate CDNA3’s 2:4 sparsity stage by stage, from the SMFMAC instructions through hipSPARSELt, and test a break-even model on shapes not used to fit it (Section VII). • We present a multi-stream benchmark for GPUs and run it on both generations under one protocol, varying streams, hardware queues, CU masks and partitions (Section VIII). • We measure FP8 and 2:4 in LLM inference, serving and training and in mixed-precision HPL, and derive guidance for library developers, system administrators and users (Sections IX and X). We will open source the benchmarks upon acceptance; we withhold them now for double-blind review. •

II. BACKGROUND : CDNA3 AND CDNA2 CDNA3 adds FP8 matrix instructions, 2:4 structured sparsity and compute partitions to CDNA2’s design, and like CDNA2, it dispatches concurrent work through hardware queues. The MI300A (Fig. 2) has 6 CDNA3 accelerator dies (XCDs) of 38 enabled CUs each, 228 in total, and its CPU cores share the GPU’s HBM3 [10].The MI250X (CDNA2) is an Open Compute Project (OCP) Accelerator Module (OAM) package of 2 graphics compute dies (GCDs) with 110 CUs each [8]. Table I lists the vendor peaks [9], [13].

(b)

Matrix core

§VIII

FP16/BF16 MFMA

§V

FP8/BF8 MFMA

§VI

2:4 sparse SMFMAC

§VII

ACEs

SIMD SIMD SIMD SIMD

Partition: SPX / TPX / CPX (SPX unless stated)

LDS (64 KiB)

new in CDNA3

on CDNA2 and CDNA3

§ = section it’s mentioned

Fig. 2. MI300A architecture. (a) The full package. (b) One compute unit, and the dispatch path from HIP streams through hardware queues to it.

FP8 matrix instructions. Both generations issue matrix work as matrix fused multiply-add (MFMA) instructions [14]. CDNA3 adds FP8 and BF8 MFMA instructions that accumulate in FP32, and the datasheet states that they run at twice the FP16 rate [13]. Their encodings are the FNUZ variants of E4M3 and E5M2, which have no infinities and no negative zero [14]. These differ from the OCP 8-bit floatingpoint (OFP8) formats [1], [2], so software must support FNUZ explicitly; as Table III shows, rocBLAS has no FP8 type. 2:4 structured sparsity. The sparse MFMA instructions (SMFMAC) execute 2:4 sparsity, in which two of every four consecutive elements are zero [10], [14], and the datasheet states that they run at twice the dense rate [13]. The saving is in arithmetic only. Pruning and compression are separate stages the application must run before the multiply, and Table III shows that hipSPARSELt does not expose every type that SMFMAC accepts. The datasheet states neither the cost of those stages nor the reuse at which the 2× pays for them, so we time each stage in Section VII. Partitions and hardware queues. Users can split the MI300A into SPX, TPX or CPX compute partitions (one, three or six partitions) [10], [15]. The MI250X has no partition modes; its GCDs enumerate as separate devices, so the GCD is the finest unit a scheduler can allocate. On both parts, hardware queues map onto asynchronous compute engines [10], [16]. Section VIII measures them with precision and sparsity varied. III. R ELATED W ORK No prior AMD work in Table II covers more than four of its seven dimensions, and none pairs FP8 or 2:4 with a generational comparison.

TABLE II C OVERAGE OF PRIOR MICROBENCHMARK STUDIES AND THIS PAPER . S H .: CONCURRENT STREAMS ; G EN .: TWO GENERATIONS . Work

Platform

Our paper

ISA LIB APP FP8 2:4 Sh. Gen. ✓ ✓ ✓ ✓

✓ ✓ ✓ –

– ✓ ✓ –

– ✓ ✓ –

– – ✓ –

– – – –

– – ✓ ✓

MI300A, MI250X ✓

✓

✓

✓

✓ ✓

✓

Schieffer [17] MI250X Hanindhito [11] MI300X/325X Jarmusch [18] B200, H200 Tee [12] MI250X/300A/X

1) Instruction and library level: Schieffer et al. [17] characterize the MI250X’s CDNA2 matrix cores with microbenchmarks and rocBLAS GEMM in FP64, FP32 and mixed FP16. We re-measure the dense path so that both sides of a generational ratio come from one protocol. Hanindhito and Patel [11] characterize the MI300X and MI325X at the instruction, library and application levels. They leave 8-bit GEMM out because of a hipBLASLt performance regression, and they measure no sparsity, sharing or CDNA2 part. Ambati and Diep [19] correct FP8 and FP16 library GEMM efficiency on the MI300X for the measured clock, as capability gain B does. They do so at one level and on one generation. NVIDIA parts have comparable studies [18], [20]–[22]. 2) Memory and partitioning: Tee et al. [12] compare the memory-latency hierarchy of MI250X, MI300A and MI300X, including the last-level cache, memory (NPS) partitioning and contention. Their generational comparison is for memory and ours is for the compute features CDNA3 adds, so we do not re-measure memory. Prior work [23]–[26] measures MI300A unified memory, copy strategies, CPU–GPU contention and inter-APU communication, and Abouelmagd et al. [15] sweep the partition modes. 3) Sharing: Otterness and Anderson study CU masking [27] and the path from HIP streams through user-mode queues to the asynchronous compute engines [28] on GCN GPUs, and show that the runtime multiplexes all streams onto a small pool of hardware queues. ECLIP [29] uses CU-masked stream pools on MI50, and RAPID-Serve [30] co-locates LLM prefill and decode with CU masking on MI300X. None of the three varies precision or sparsity or compares two generations. We do both, with the queue count as a controlled variable. 4) Sparsity: Many studies cover 2:4 sparsity on NVIDIA GPUs [3], [31]–[34]. On AMD, Dickmann et al. [35] compare a 2:4-sparse LLM on MI300A with a larger dense model, which confounds sparsity with model size, and do not report whether the sparse hardware path runs. To our knowledge, no AMD study separates pruning and compression from the multiply or reports a break-even point. We time each hipSPARSELt stage, test a break-even model on held-out shapes, and run the hardware 2:4 path under sharing. IV. M ETHODOLOGY A. Hardware setup We run on two systems with nodes for MI300A and MI250X to compare package against package. The MI300A

runs in SPX/NPS1 mode with a power cap of 550 W per APU socket, and the MI250X is capped at 560 W per OAM. A one-GCD run leaves the sibling GCD idle, and a package run uses both. We compare packages because the GCDs share one power cap and only one reports the package’s power; one GCD is a secondary comparison. B. Experiment methodology We use one benchmark per level, from single matrix instructions to whole applications, and Sections V–IX report their results in this order. All run on ROCm 7.2.0, the LLM benchmarks use PyTorch 2.7.0+rocm6.3 (HIP 6.3.42131), and the artifact lists the library, driver and OS versions. 1) Matrix-instruction microbenchmarks: Our microbenchmarks time single matrix instructions with no library involved. Each issues one MFMA or SMFMAC compiler builtin in a loop whose operands stay in registers. Throughput variants run independent accumulator chains in every wavefront across the device; latency variants run one wavefront through a chain of dependent MFMAs. Each wavefront reads its shader clock and a constant-rate counter around the loop, which yields cycles per instruction. We build each microbenchmark with the default code generation and with a form that keeps accumulators in the matrix registers and inspect the disassembly. 2) Library benchmarks: We time each GEMM in hipBLASLt, rocBLAS and hipSPARSELt with its fastest kernel, so the results reflect the library and not its heuristic. The shapes are N 3 for N from 1024 to 16384 in powers of two, an MLP shape (M ×N ×K = 4096×16384×17408) and a thin-K shape (8192×8192×1024). Every GEMM uses the NN layout and FP32 accumulation, FP8 GEMMs write FP16, and throughput is 2M N K over the median time. We compare two libraries only when the input, output and accumulator datatypes all match. For hipBLASLt we time every algorithm that getAllAlgos returns for those datatypes and keep the fastest; rocBLAS runs rocblas_gemm_ex with its standard algorithm. We time hipSPARSELt’s prune, compress, plan and matmul stages separately, and run the sparse matmul and the matching dense GEMM at their best algorithms. 3) Multi-stream benchmark: The multi-stream benchmark runs several GEMM streams on one device at once and varies how many hardware queues they share. Each HIP stream runs the same GEMM once per round, always with the same kernel, chosen before the run. We vary the stream count over 1, 2, 4, 8, start the streams together with a flag set on the GPU, and shuffle their launch order with a recorded seed. HIP maps streams onto a default pool of four hardware queues [28], so from four streams up, streams share queues. We run every configuration with this default and with 16 queues, where no two streams share one. To divide the MI300A among streams that do the same total work, we also limit each stream to its own CUs with CU masks, or split the device into three or six separate devices with TPX and CPX partitions. 4) Application benchmarks: The application benchmarks measure FP8 and 2:4 in one LLM, in inference, serving and training, and the FP16 matrix path in mixed-precision HPL.

LLM. We run Llama-3.1-8B-Instruct with 512-token prompts on a plain PyTorch/transformers harness that keeps every kernel visible, and time each phase (prefill and decode) and each layer. Attention uses PyTorch’s MATH backend, checked against an FP32 reference before each run; fusedattention (FLASH) runs are reported beside it. We time 2:4 layers in hipSPARSELt against a tuned dense baseline at every linear shape, on the MI300A only, and measure WikiText-2 perplexity over 262,143 tokens with the same masks on both machines, without fine-tuning. We serve the model in three ways: with vLLM [36], which fuses activation quantization (FP16 and FP8 on the MI300A, FP16 on the MI250X GCD); as several processes with their own replicas on one device; and as 6 processes in CPX partitions against the same processes in SPX. Serving reports throughput, time to first token and time per output token over a fixed request trace and a mixed-length one. For training, we measure the time per step of a GPTstyle model with 1.48 billion parameters at 16,384 tokens per step. The MI300A trains it in BF16 and in FP8, with torchao 0.18.0 converting the linear layers to FP8 with dynamic pertensor scaling; the MI250X GCD and package train it in BF16. Every run starts from the same weights and data order, and we compare the losses step by step. The training runs use an MI300A node capped at 760 W. HPC. We run mixed-precision HPL [4], [5] in AMD’s unmodified rocHPL-MxP 2.1.0 [37]. Each solve of a dense system Ax = b factors A with a rocBLAS FP16 GEMM in its LU trailing update (K = 2, 560) and refines the result in FP64; it passes only if its scaled FP64 residual falls below 16. Each device runs one MPI rank, and the package one solve per GCD. We time 32 solves per size with the benchmark’s host clock and discard the first 2. We also run a STREAM Triad and rocBLAS DGEMM. C. Capability gain Capability gain measures how much of a capability’s instruction-level speedup a level keeps. For a capability (FP8 over FP16, or 2:4 sparse over dense) at level ℓ, it is Gℓ = Sℓ / SISA ,

(1)

where Sℓ is the measured speedup at level ℓ and SISA is the instruction-level speedup. Gℓ = 1 means the level keeps the whole speedup, and Gℓ > 1 means it is faster than the instructions predict. We take SISA as the architectural ratio, 2.00, because the measured FP8/FP16 clock ratio varies with the operands (Section VI). Across generations, capability gain B divides a generational speedup by the architectural FLOP-per-CU-per-cycle ratio, the CU-count ratio and the measured clock ratio. D. Measurements Every result is the median of repeated runs without profiling. HIP events time each measured region, which ends with a stream synchronization, after at least 10 warmup iterations and over at least 30 stored repeats, and we report the interquartile range with each median. We randomize the configuration order

TABLE III H OW FP8 AND 2:4 ARE REACHED AT EACH LEVEL ON THE MI300A, AND WHERE THE MI250X PATH STOPS . FP8 GEMM MI300A (CDNA3) ISA v_mfma_f32_32x32x16_ fp8_fp8 Library hipBLASLt with HIP_R_8F_ E4M3_FNUZ; rocBLAS has no FP8 type PyTorch torch._scaled_mm, float8_e4m3fnuz; OCP e4m3fn returns NaN MI250X (CDNA2): where the path stops ISA compile error (no fp8-insts) Library hipBLASLt: no FP8 algorithm PyTorch _scaled_mm is MI300+ only

2:4 sparse GEMM v_smfmac_f32_32x32x16_ f16 hipSPARSELt prune → compress → plan → matmul, for FP16, BF16 and INT8; FP8 error –

backend cannot select SMFMAC hipSPARSELt: errors for every type –

with a recorded seed. We use hardware counters only for diagnosis and for capability gain B, because profiled runs serialize dispatches and inflate clocks for short kernels. We take B’s clock from each XCD’s hardware counter of busy cycles, check it against a second such counter in the same pass, and report B only at shapes where the two agree within 2.0% on both GPUs. We normalize throughput only by measured peaks, the highest sustained throughput in a dedicated sweep, which Table IV lists. Library throughput varies by up to 4.0% between two MI300A devices and by up to 6.71% between the two GCDs of one MI250X, so we treat any smaller difference as noise. The values in the input matrices change GEMM throughput and the MI250X’s delivered clock (Section VI), so we fix them: every measured peak multiplies the same random input values (canonical_rand), generated from a fixed randomnumber seed. We set the MI250X’s power-management (DPM) level to manual and measure its delivered clock inside the kernel with extra wavefronts that read the clock, not from the driver. The MI300A runs under default power management. V. I NSTRUCTION L EVEL : M ATRIX C ORES RQ1 starts at the instructions, and there the speedup survives: CDNA3 delivers AMD’s advertised speedups per issue cycle, and later sections measure against them. We rescale the rates by each GPU’s measured shader clock. The rescaled rates match the datasheet peaks [8], [9], [13] to within 0.02% on CDNA3 and 0.05% on CDNA2, so any shortfall from the vendor peak lies in the clock. Table V shows that one MI300A runs FP16 MFMA 4.30×–4.69× and FP8 MFMA 8.50×–9.11× faster than one MI250X GCD runs FP16. 1) Tile size: AMD’s claim that CDNA3 doubles CDNA2’s FP16 rate per CU per cycle [10] holds only for 32×32 tiles. CDNA3 issues at 2.00× CDNA2’s rate for 32×32 tiles but at only 1.78× for 16×16 tiles, so kernels that use 16×16 tiles get less than the advertised doubling. FP8 has no such limit: one FP8 MFMA instruction performs 2.00× the FLOPs of an FP16 MFMA instruction in the same number of issue cycles

TABLE IV M EASURED GEMM PEAKS AND HBM BANDWIDTH OF THE MI300A, ONE MI250X GCD AND THE MI250X PACKAGE . B EST MEASURED RATE PER LIBRARY; HBM: DEVICE - TO - DEVICE COPY KERNEL ; per CU·clk: THE GCD’ S RATIO PER CU AND PEAK CLOCK . MI300A Prec.

Library

TFLOP/s

Device

MI250X, one GCD

% peak

TFLOP/s

228 CUs, 2.1 GHz

% peak

vs MI300A

MI250X package per CU·clk

TFLOP/s

110 CUs, 1.7 GHz

FP64 FP64 FP32 FP32 FP16 FP16 BF16 BF16 FP8

rocBLAS hipBLASLt rocBLAS hipBLASLt rocBLAS hipBLASLt rocBLAS hipBLASLt hipBLASLt

69.4 16.3 91.7 94.1 566 565 590 590 1,080

57% 13% 75% 77% 58% 58% 60% 60% 55%

26.1 6.3 40.4 40.1 128 142 139 153

55% 13% 84% 84% 67% 74% 73% 80%

HBM

TB/s

3.24

61%

1.17

71%

% peak

vs MI300A

220 CUs, 1.7 GHz

0.377× 0.386× 0.441× 0.427× 0.226× 0.251× 0.236× 0.260×

0.965× 0.989× 1.13× 1.09× 0.578× 0.641× 0.603× 0.666×

50.7 12.7 78.0 75.3 244 273 262 287

53% 13% 81% 79% 64% 71% 68% 75% –

0.731× 0.778× 0.850× 0.800× 0.432× 0.483× 0.444× 0.486×

0.360×

–

2.34

71%

0.721×

–

TABLE V M ATRIX - INSTRUCTION THROUGHPUT, ISSUE CYCLES AND DEPENDENT- ISSUE LATENCY ON CDNA3 AND CDNA2; CDNA3: THE WHOLE MI300A; CDNA2: ONE GCD. FLOP/CU/ CYC AT THE PEAK CLOCK ; ISSUE CYCLES AT ONE WAVE PER SIMD; LATENCY OVER A 32- DEEP DEPENDENT CHAIN . Arch.

Prec.

Tile

TFLOP/s

% ISA peak

% vendor

FLOP/CU/cyc

Issue cyc

Latency cyc

vs CDNA2 FP16

CDNA3 CDNA3

FP8 FP8

16x16x32 32x32x16

1,520 1,596

95.3% 100%

77.5% 81.4%

3,175 3,334

18.0 32.0

16.9 32.4

8.50× 9.11×

CDNA3 CDNA3

FP16 FP16

16x16x16 32x32x8

769 822

93.6% 100%

78.5% 83.9%

1,607 1,718

18.0 32.0

16.9 32.4

4.30× 4.69×

CDNA3 CDNA3

BF16 BF16

16x16x16 32x32x8

778 816

95.3% 100%

79.4% 83.3%

1,626 1,705

18.0 32.0

16.9 32.4

4.35× 4.66×

CDNA2 CDNA2

FP16 FP16

16x16x16 32x32x8

179 175

100% 98.0%

93.4% 91.5%

956 937

32.0 64.0

32.5 64.1

1.00× 1.00×

CDNA2 CDNA2

BF16 BF16

16x16x16 32x32x8

179 175

100% 97.9%

93.4% 91.4%

956 936

32.0 64.0

32.5 64.1

1.00× 0.999×

for every tile size (12 matched pairs), so FP8’s instructionlevel speedup does not depend on which tile the library picks. 2) Code generation: The compiler sets CDNA3’s measured 16×16 instruction peak, so we report the best build. The default build wraps CDNA3’s MFMAs in extra instructions that move accumulators, and the register-form build, which keeps accumulators in the matrix registers, runs 1.41–1.43× faster at 16×16. 3) Latency and occupancy: FP8 and FP16 MFMA instructions of the same tile have the same latency, so FP8’s 2× comes from doing twice the work per instruction, not from finishing sooner. On both generations, a dependent MFMA waits about one issue interval (the cycles one instruction occupies the matrix pipe): 16.9 and 32.4 cycles for 16×16 and 32×32 tiles on CDNA3. A chain of dependent MFMAs therefore stalls little. CDNA3 reaches its measured FP8 and FP16 peaks with one wavefront per SIMD.

in FP16 on the MI250X package, close to what the hardware predicts. Table VI lists each library path; CDNA2’s baseline is its FP16 path at the same shape, datatypes and input values, with BF16 as an alternative.

VI. FP8 T HROUGH THE L IBRARY AND ACROSS A G ENERATION The answer to RQ1 is mostly yes: hipBLASLt keeps most of FP8’s instruction-level 2.00× except in small GEMMs, and the library’s kernel choice, not the hardware, causes the loss. An 81923 GEMM runs 3.95× faster in FP8 on the MI300A than

2) Cause of the loss: hipBLASLt’s kernel choice causes the loss at the MLP shape, the one shape we traced. Hardware counters show that hipBLASLt picks an FP8 tile twice as deep in K as the FP16 one (512x144x64 against 512x144x32) and re-reads its operands 2.81× as often. FP8 therefore moves 1.39× FP16’s bytes per FLOP, which cancels the benefit of its

1) Library capability gain: hipBLASLt delivers most of FP8’s speedup in large GEMMs and loses much of it in small ones. Fig. 1(a) shows that the FP8 library capability gain ranges from 0.580 at 10243 to 1.01. The clock does not explain the loss: at the 3 shapes where we measured the FP8/FP16 clock ratio, dividing by it as well changes the capability gain by at most 3.9%. The capability gain exceeds 1.0 at 1 of 7 shapes (20483 ) because the FP16 kernel is weak. hipBLASLt’s own heuristic does not always pick the best kernel: its top choice costs up to 25% against the fastest kernel the library offers once the process has used the library, and in a new process it can be 35.5× slower than the best. The heuristic costs the MI250X less than its GCD spread.

40963

thin-K

81923

MLP

163843

– 125 119 117 121

89.6 93.1 76.0 85.2 –

103 146 133

80.4 95.4 94.0 –

126 150 135

98.2 97.8 95.7 –

119 152 139

93.0 99.3 98.0 –

137 119

98.4 93.3

halved operand width. The same counters rule out conversion traffic, spilling, local data share (LDS), occupancy and clock. 3) Generational speedup: Upgrading to the MI300A and adopting FP8 together make an 81923 GEMM 3.95× faster, package against package. We split this speedup into the FP16 generational speedup (CDNA3-FP16 over CDNA2-FP16 at the same configuration) and the FP8 library capability gain above. Fig. 3 compares their product with the measured end-to-end speedup. The FP16 generational speedup is 1.20–2.15×. 4) Per-CU·clock bound: The FP16 generational speedup is close to what CDNA3’s hardware predicts: it reaches 0.870– 0.980 of the per-CU·clock bound, so CU count, clock and the doubled matrix rate explain almost all of it, and FP8 end to end reaches 0.831–0.872 of its bound. The bound has to use measured clocks. The measured clock ratio, 1.02–1.05×, is much closer to one than the peak ratio 1.24×, and operand values move the clock: zero operands run up to 44.5% faster on the MI300A, and the MI250X driver hides a drop in its delivered clock. Fig. 3 marks the shapes where both clocks were measured; neither the one-GCD scope nor CDNA2’s faster BF16 baseline (dotted) changes these conclusions. VII. S TRUCTURED S PARSITY AND I TS B REAK -E VEN For RQ2, 2:4 loses its speedup at two boundaries: the instruction clock costs part of it, and hipSPARSELt, which keeps only 0.376–0.554 of the 2× on the MI300A, loses the rest in

CDNA3 FP8 CDNA3 FP16

16K3

8K3

thin-K

77.8 85.2 86.6

no speedup MLP

– 99.5 131 123

1

4K3

46.7 66.4 69.2

2

2K3

– 59.7 102 98.0

2.61× 2.41× 2.79× 1.49× 2.91× 2.67× 5.41× 2.16× 4.31× 4.29× 7.63× 2.66× 2.48× 3.95× 3.82× 6.25× 3.44× 4.37× 4.19× 7.89× 3.43× 4.37× 4.25× 7.62× 3.66× 4.33× 3.77× 7.49× 4.28× 4.11×

1K3

35.3 40.3

16K3

54.1 57.1

8K3

25.3 24.3 15.2 15.0 48.4 46.2 50.5 37.4 89.6 93.2 89.2 55.3 53.7 80.7 81.4 71.8 73.3 98.5 98.6 100 80.6 100 102 98.4 87.7 102 92.5 99.0 100 101

4

thin-K

20483

149 138 160 85.1 285 261 530 212 529 527 936 326 304 476 460 753 415 581 557 1,050 456 591 576 1,032 496 601 523 1,039 594 570

Lt FP16

MLP

Lt BF16 Lt FP16 Lt FP8 rB FP16 Lt BF16 Lt FP16 Lt FP8 rB FP16 Lt BF16 Lt FP16 Lt FP8 rB BF16 rB FP16 Lt BF16 Lt FP16 Lt FP8 rB FP16 Lt BF16 Lt FP16 Lt FP8 rB FP16 Lt BF16 Lt FP16 Lt FP8 rB FP16 Lt BF16 Lt FP16 Lt FP8 rB BF16 rB FP16

vs MI250X

4K3

10243

MI250X GCD

TFLOP/s % meas. TFLOP/s % meas.

(b) over the MI250X package

8

2K3

Path

(a) over one MI250X GCD

1K3

MI300A Shape

speedup over CDNA2 FP16 (×)

TABLE VI L IBRARY GEMM THROUGHPUT AND SHARE OF EACH PART ’ S MEASURED PEAK (% MEAS .), MI300A VS ONE GCD. LT: HIP BLASLT, R B: ROC BLAS.

clock bound vs faster 16-bit

Fig. 3. CDNA3 FP8 and FP16 GEMM speedup over CDNA2 FP16 on the MI250X package, per shape. Solid: against CDNA2 FP16; dotted, hollow: against its faster 16-bit path. Ticks: the per-CU·clock bound, where both clocks were measured.

keeping the matrix units supplied with data. A cost model built from each stage predicts pruning within 2.9% but not the break-even point, and hipSPARSELt’s kernel choice adds an uncertainty of up to 6.0–8.8% to every sparse-versus-dense ratio. The SMFMAC FP8 instruction runs, but hipSPARSELt supports 2:4 only for FP16, BF16 and INT8 [38], as Table III shows. 1) Instruction level: 2:4 sparsity loses part of its expected 2× even at the instruction level, without any library. One SMFMAC instruction performs twice the multiply-accumulates of the dense MFMA instruction with the same tile and precision in the same number of issue cycles (1.99–2.04 over 10 pairs), so the 2× holds per cycle. The sparse instructions, however, run the MI300A at 0.786–0.908 of the dense clock, so the delivered speedup is only 1.57–1.84. We did not measure power in these runs, so we cannot attribute the lower clock; powerlimited dynamic voltage and frequency scaling (DVFS) under the APU’s single power cap is one candidate. hipSPARSELt then loses more. 2) Library speedup: hipSPARSELt’s sparse matmul is at best slightly faster than dense, and usually slower. Fig. 1(b) shows its capability gain (Eq. (1)) over the 25 shapes we use to fit the model below: it ranges from 0.376 to 0.554, and the sparse matmul wins at only 8 of 25 of them. The best case, an LLM projection-layer shape held out from the fit, reaches a median speedup of 1.21–1.23 across independent processes. 3) Cause of the loss: The sparse kernels lose their speedup in keeping the matrix units supplied with data, not in the matrix units themselves or in the clock. Counters show that SMFMAC delivers 1.00 of its 2× per cycle the matrix unit is busy (0.93–1.03 over 27 of 36 shapes). The sparse kernel, however, keeps the matrix pipe busy only 0.46 as often as the dense hipBLASLt kernel for the same GEMM (0.39–0.60) and issues 1.92× the LDS instructions per FLOP. It runs at 1.03–1.12× the dense kernel’s clock at the 5 shapes where we measured both, so the clock does not cause the loss. The

10

7

10

6

10

5

M ×K

Prune

Compress

Sparse/

µs

µs

error

dense

+0.9% 117 +0.0% 342 -0.6% 396 +0.3% 392 +0.4% 111 +2.9% 177 +1.4% 2,745

+40.9% +3.2% +0.2% -0.9% +51.2% +20.5% +16.6%

0.912 0.886 0.917 0.840 0.938 1.05 1.00

3072 × 3072 147 6144 × 6144 584 4096 × 11008 693 4096 × 11008 (BF16) 699 8192 × 1024 130 1024 × 16384 267 16384 × 16384 4,213

error

break-even N *

TABLE VII HIP SPARSELT PRUNE AND COMPRESS TIMES ON THE MI300A AT SHAPES FROM THE MODEL FIT, AND THE MODEL’ S ERROR ( MEASURED / MODEL −1). S PARSE / DENSE : MATMUL TIME RATIO , MEDIAN OVER N .

c + γ/M K , 2 (1/Rde − 1/Rsp )

H5

10

4

10

3

H4 H1

H2

H3b H3

FP16: the model gives no break-even

10

7

10

8

M×K

instruction-level loss, by contrast, comes from the lower clock, so the library loses its speedup for a different reason than the instructions do. 4) Break-even model: A cost model built from each stage predicts the pruning cost, but not the break-even point. Sparsity pays off only when the matmul runs on enough columns N to cover the cost of pruning and compressing the weights once. Pruning and compression both scale with the weight size: Tprune = a M K and Tcomp = b M K, with c = a + b. Planning costs a constant γ, and each matmul runs at its measured peak rate, T = 2M N K/R, with Rde for dense and Rsp for sparse. The break-even width N ∗ is where the one-shot sparse total, which pays every stage once, equals the dense matmul: N∗ =

model N*, BF16 measured, FP16 measured, BF16 held-out N grid

(2)

which is independent of M and K once γ/M K is small.We fit the constants on these 25 shapes and test them, unchanged, on 7 held-out shapes, which Table VII lists; a cell is one shape at one value of N . The model times pruning within 2.9% at every held-out shape but misses compression by up to 51.2% at the smallest M K, which points to a fixed compression cost it lacks. It also misses how the break-even depends on shape: Fig. 4 shows that the measured FP16 crossings span 4.78× in N ∗ , which Eq. (2) cannot produce, and the model places the BF16 N ∗ 8.7× above the measured crossing, because its constants assume a 1.5% per-FLOP advantage for sparse where the measured shape delivers 17–21%. The FP16 ratio of the measured peaks, Rsp /Rde = 0.983, predicts no break-even at all, yet the sparse matmul alone is faster at 18 of 30 held-out cells, because the two libraries reach their measured peaks at different shapes. Users therefore have to measure a sparse-versus-dense ratio per shape. hipSPARSELt also picks different kernels in different processes, at 5 of 6 cells even with the node and binary held equal. We call the largest spread of the end-to-end ratio across such processes the selection floor, 6.0–8.8%; it makes every sparse-versus-dense ratio uncertain by that much, so in 4–8 of 18 held-out FP16 cells where sparse measured faster, the margin is too small to call.

Fig. 4. Measured 2:4 break-even crossings N ∗ against the base model, by M K. Circles: FP16; squares: BF16; line: the model’s BF16 N ∗ ; bars: heldout N grids, with no crossing.

VIII. C APABILITY G AIN U NDER S HARING For RQ3, the boundary is the hardware-queue pool: it, not precision, sets fairness on a shared device: with enough queues the pool treats every precision equally, and the default four queues trade fairness for throughput. FP8’s speedup over FP16 grows when streams share queues, and the growth tracks a larger clock drop for FP16. 2:4 sparsity also tends to gain under sharing, and hard partitions isolate streams better than CU masks. 1) Metrics: Three metrics compare each stream’s shared wall time with its solo time. Jain’s index J [39] measures how evenly sharing slows the streams; it is 1 when all slow down equally, and we compare it only across streams that run the same GEMM. System throughput (STP) sums each stream’s solo-over-shared time, so it reaches n for n perfectly overlapped streams and is 1 when sharing is no faster than running them one after another. Average normalized turnaround time (ANTT) is the mean shared-over-solo time, that is, the average slowdown. 2) Fairness: The number of hardware queues, not the numeric precision, decides how fairly concurrent streams share the device, and changing the queue count trades fairness for throughput. Table VIII shows that with 16 queues, from four streams up, J is at or near one for every precision on both GPUs, ANTT grows roughly as n, and STP stays near one. At n = 2 the two MI300A streams take turns, one at its solo speed, so J is low and STP far below two. With the default 4 queues, both sizes we ran with eight streams lose fairness and gain throughput: on the MI300A, J falls to 0.875–0.942 at 40963 and 0.927–0.938 at 81923 , while FP16 STP rises from 0.96 to 2.01 and from 0.63 to 1.12. The MI250X GCD shows the same trade. 3) FP8: FP8’s speedup over FP16 grows when streams share hardware queues. The MI300A’s FP16-over-FP8 walltime ratio at 81923 with 16 queues rises from 1.92× alone to 2.46× at eight streams (+28.6%), and the full ranges, [1.91, 1.92] and [2.45, 2.50], do not overlap. Capability gain (Eq. (1)) rises from 0.958 alone to 1.29 at two streams and is still

TABLE VIII FAIRNESS AND THROUGHPUT OF CONCURRENT STREAMS BY PRECISION AND STREAM COUNT n AT 40963 ; J , STP AND ANTT WITH 16 HARDWARE QUEUES , J4 WITH THE DEFAULT 4. J : JAIN ’ S FAIRNESS INDEX ; STP: SYSTEM THROUGHPUT; ANTT: AVERAGE NORMALIZED TURNAROUND TIME . MI300A J4 STP ANTT

J

J4 STP ANTT

Harness

– – – –

– – – –

vLLM

n

dense FP8

1 1.00 1.00 1.00 2 0.910 0.913 1.50 4 0.996 0.960 0.98 8 1.00 0.875 1.14

1.00 1.49 4.12 7.02

– – – –

– – – –

1 1.00 1.00 1.00 2 0.907 1.00 1.51 4 0.997 0.922 0.98 8 1.00 0.875 0.96

1.00 1.00 1.00 1.00 1.48 0.993 0.997 1.08 4.12 1.00 0.939 1.01 8.33 0.999 0.924 0.88

1.00 1.86 3.97 9.11

2:4 sparse FP16 1 1.00 1.00 1.00 2 0.988 0.992 1.31 4 1.00 0.886 1.06 8 1.00 0.942 1.01 dense BF16

1 not measured 2 not measured 4 not measured 8 not measured

1.00 1.56 3.79 7.93

FP8 Workload Part

MI250X GCD

Streams

dense FP16

J

TABLE IX L LAMA -3.1-8B SPEEDUPS ON THE MI300A. FP8: OVER FP16 (BF16 FOR TRAINING ), STATIC ( S ) OR DYNAMIC ( D ) SCALE . G EN .: FP16 (BF16) OVER ONE MI250X GCD AND THE PACKAGE .

– – – –

– – – –

– – – –

– – – –

1.00 1.00 1.00 0.999 0.999 1.04 1.00 0.920 1.03 1.00 0.943 1.04

1.00 1.93 3.90 7.66

Training

prefill decode fixed trace mixed trace step

Gen.

s

d

1.01–1.35 0.88–1.03 1.17–1.38 1.16–1.32 –

0.79–1.10 0.59–0.88 1.13–1.44 1.13–1.42 1.15

/GCD

/pkg.

1.97–2.94 – 1.23–2.20 – 2.26–2.96 1.36–1.73 2.27–2.82 1.39–1.65 3.31 1.94

IX. A PPLICATION C ASE S TUDIES To answer RQ4, how much of FP8’s speedup reaches applications depends on fusion. FP8 speeds up the LLM only where its conversion is fused (vLLM serving, 1.13–1.44×), not in our unfused harness, and a training step by only 1.15×; 2:4 speeds up no layer; and on the MI250X the FP16 LU of rocHPL-MxP does no work, so that benchmark yields no generational ratio. A. FP8 in LLM inference, serving and training

1.23 at eight, so something other than arithmetic now drives the speedup. Queue sharing causes the growth: it takes the ratio from 1.85× alone to 2.48× at six streams, while CPX partitions (one stream per XCD) leave it flat, 1.73× to 1.71×. The extra speedup tracks the clock. Power management under the APU’s single power cap throttles FP16 harder than FP8, so the FP8/FP16 clock ratio grows from 1.07 alone to 1.36 at eight streams, which explains 95% of the growth in capability gain; a second node repeats this. The attribution is consistent with DVFS but, with 2 repeats per configuration on two nodes, does not prove it. The power cap does not explain everything: with eight streams at 20483 the socket stays below its cap, yet FP16 still clocks lower than FP8. 4) 2:4 sparsity: Sharing tends to favour 2:4 sparsity, but only the direction is reliable. Sparse runs about as fast as dense FP16 alone at 81923 and 1.39× faster at eight streams, and sharing favours sparse at 4 of 6 shapes. Swapping only the sparse algorithm moves the eight-stream ratio by +8.6%, and the selection floor of Section VII is 6.0–8.8%, so we do not report the size of the effect. 5) Dividing the GPU between streams: Hard partitions isolate streams best, and CU masks worst. Partitions deliver 1.31× and 1.29× the total throughput of queue sharing for three and six FP16 streams, while CU masks deliver only 0.78× and 0.70×. A mask cannot confine a stream to fewer than 6 XCDs and fails silently, so masked streams still share every L2; we hypothesize that partitions win by giving each stream whole XCDs.An administrator who shares an MI300A between jobs should therefore partition it rather than use CU masks.

FP8 speeds up the LLM only where its conversion to FP8 is fused. Table IX lists FP8’s speedup and the MI300A’s generational speedup for Llama-3.1-8B-Instruct in each mode; our plain PyTorch harness is slower than a production engine, but both GPUs run the same code, so the ratios are fair. 1) Unfused inference: Prefill exercises the matrix engines, and the MI300A’s FP16 prefill is 1.97–2.94× faster than one GCD’s. Decode is bound by the CPU launching kernels, not by the GPU: at batch 1 a step takes 5.8× the time needed to stream the weights once. Fused attention speeds an FP16 request at batch 32 by 1.42× on the MI300A and 1.22× on a GCD, with the same first token. FP8 runs on the MI300A but gains little. PyTorch’s OCP type float8_e4m3fn returns non-finite results on the MI300A, whose FP8 instructions use the FNUZ encoding (Section II). With a static scale, calibrated in advance, FP8 is 1.12× faster at prefill and 0.889× at decode, a loss, and a dynamic scale, computed per call, lowers decode to 0.594×. The FP8 kernel alone is 1.55× faster at the model’s layer shapes, but the whole layer delivers 0.862×: the harness converts activations to FP8 in separate kernels, at a median 27.9 µs per call, and at 10 of 50 cells the conversion alone costs more than the FP16 matmul it replaces. This cost is an upper bound on what a fused engine pays. 2) Fused serving (vLLM): A serving engine fuses activation quantization into the preceding kernel and serves many requests at once. We serve the model with vLLM at several concurrency levels (Section IV-B). Here Sℓ in Eq. (1) is the engine’s FP8-over-FP16 throughput speedup. Fusing the quantization makes FP8 pay off at every load we serve. FP8 increases the MI300A’s output throughput by 1.13–1.44× with a dynamic scale and 1.16–1.38× with a calibrated static scale, capability gains of 0.564–0.719 and

(a) loss per step

(b) median step 2,129

loss

9 MI300A FP8

8 7 MI300A BF16 MI250X BF16

50 100 training step

step time (ms)

2000 1,245

1000

643

557

0 GCD pkg. BF16 FP8 BF16 BF16 MI250X MI300A

Fig. 5. (a) Loss per timed step, BF16 and FP8 on the MI300A and BF16 on one MI250X GCD (dotted). (b) Median step time (bars) and each repeat’s median (dots).

0.580–0.689, whereas the unfused harness above loses at decode. In FP16 the MI300A serves 2.26–2.96× one MI250X GCD’s throughput and 1.36–1.73× the package’s, which we run with tensor parallelism over both GCDs. At 64 requests FP8 cuts the median time to first token from 119 to 79.5 ms. 3) Training: Inference loses FP8 at small, launch-bound GEMMs; training runs large GEMMs, where hipBLASLt keeps nearly all of FP8’s instruction-level speedup (1.01 at 163843 , Section VI). Section IV-B describes the training setup. Fig. 5 shows that FP8 shortens a training step far less than its GEMM speedup predicts and that its loss stops falling. An FP8 step is 1.15× faster than BF16 and launches 1.65× as many kernels, the same conversion cost as in inference. Its loss lies up to 36% above BF16’s, so the speedup is per step, not per unit of training progress; we did not tune the FP8 recipe. In BF16 the MI300A step is 3.31× faster than one GCD’s and 1.94× faster than the package’s, and fused attention speeds it by a further 1.08×. 4) FP8 accuracy: FP8 costs little accuracy. We measure WikiText-2 perplexity, roughly the number of equally likely next tokens the model chooses among, so lower is better. Dense FP16 scores 6.39 on the MI300A and 6.39 on one GCD, 0.00570% apart, and fused attention changes the GCD’s by −0.021% and the MI300A’s by −0.0001%. The per-tensor FP8 layers above increase perplexity by 0.85% [0.79, 0.90] with a dynamic scale and 1.02% [0.96, 1.09] with a static one. An FP8 inference or training stack therefore has to fuse the conversion before FP8’s GEMM speedup can show. B. 2:4 sparsity in LLMs 2:4 speeds up none of the model’s layers and costs a lot of accuracy without fine-tuning. 1) Sparse transformer layers: 2:4 does not speed up the model’s linear layers on the MI300A, the only GPU with a 2:4 path, as Table III shows. The median dense/sparse ratio with weights compressed once and reused is 1.06×, inside the selection floor of Section VII (6.0–8.8%): at that floor, 10–15 of 43 cells are indistinguishable, 9–12 of 43 dense-faster and 19–21 of 43 sparse-faster. Dense wins at 43 of 43 cells when each call pays for compression.

2) Accuracy of 2:4 pruning: 2:4 costs a lot of accuracy. Magnitude pruning of the 224 decoder linear layers increases perplexity from 6.39 to 231 (×36.2), and hipSPARSELt’s own pruning pattern over all layers, the pattern timed above, to 10,418 (×1630), which is close to random text. Usable 2:4 therefore needs an accuracy-aware pruner and fine-tuning, which we do not evaluate. The loss comes from pruning, not from hipSPARSELt’s arithmetic, whose error matches torch’s own FP16 product. Until hipSPARSELt keeps more of its 2× and pruning keeps accuracy, dense remains the better choice for this model. C. LLM serving under sharing Whole models gain less from sharing a device than the GEMM streams of Section VIII. 1) Multi-process serving: Several processes serving their own FP16 replicas overlap little: at n = 2, J is 1.00 and STP is 1.07 on the MI300A and 1.13 on the GCD, out of 2, and the MI300A runs four processes only slightly better than serially (STP 1.08 with 4 queues, 1.11 with 16). Unlike the GEMM streams, sharing does not turn FP8 into a speedup: an FP8 process alone takes 1.17× the FP16 process’s time, and four FP8 processes take 1.19× the time four FP16 processes take to finish with 4 queues and 1.16× with 16. 2) Partitioned serving: At the GEMM level, hard partitions beat queue sharing and CU masks (Section VIII). We compare CPX against SPX with whole models by aggregate throughput, per-process latency and J (Section IV-B). For whole models the result reverses (Table XI): 6 CPX partitions deliver 0.78–0.96× the throughput of the same processes sharing SPX, although a request on one partition, with a sixth of the CUs, takes only 2.6× as long as on the whole device; one process per MI250X GCD delivers 1.76× that of both on one GCD. The GEMM-level advantage of FP8 under sharing does not carry over to whole models served side by side. D. HPC workloads rocHPL-MxP’s LU trailing update is the rocBLAS FP16 matrix-core GEMM of our measured peaks, so the question is how much of the FP16 measured peak reaches time to solution (Section IV-B describes the setup). rocHPL-MxP has no FP8 or BF16 path, and we port none. Table XII shows that every timed solve passed. The MI300A’s rate rises with N and reaches Φ = 0.608 of the rocBLAS FP16 measured peak at the largest size. The MI300A solves 2.35× faster than one GCD at the largest common size, and the package delivers 2.00× one GCD’s throughput. On the MI250X the FP16 LU does no work. The benchmark scales its generated matrix so that at least 99.9% of the FP16 operands fall below FP16’s smallest normal, and the MI250X’s FP16 matrix instructions flush subnormal inputs to zero [40]: a diagnostic build shows that the trailing updates receive only zeros, and the solve still passes because FP64 refinement, preconditioned by the FP32 LU of the diagonal blocks, needs 4–5 steps instead of 3. On the MI300A the

TABLE X MI300A PERFORMANCE SUMMARY ACROSS WORKLOADS . Workload FP16 MFMA FP8 MFMA 2:4 SMFMAC FP16 GEMM, hipBLASLt FP8 GEMM, hipBLASLt 2:4 GEMM, hipSPARSELt Shared GEMM, 8 streams LLM prefill, batch 32 vLLM, 64 requests LLM training step LLM layers, 2:4 Multi-process, n=2 Partitioned serving, 6 proc. rocHPL-MxP STREAM Triad DGEMM, rocBLAS

Metric MI300A FP8 or 2:4 vs. MI250X Limiting factor TFLOP/s 769–822 – 4.30×–4.69× 2× per CU·cycle only for 32×32 TFLOP/s 1,520–1,596 2.00× per cycle – none vs. dense – 1.57–1.84× – lower clock measured peak 565 TFLOP/s – 1.20–2.15× (pkg.) CU count and clock measured peak 1,080 TFLOP/s gain 0.580–1.01 3.95× (pkg.) kernel choice capability gain – 0.376–0.554 – feeding the matrix units FP8 over FP16 – 2.46× (1.92× alone) – queue pool and clock time (ms) 1,051 layer 0.862× 2.90× unfused conversion output tok/s 2,887 1.13–1.44× 2.96×, 1.73× none (fused) time (ms), BF16 643 1.15× 3.31×, 1.94× conversion kernels dense vs. sparse – dense wins 43 of 43 – library and pruning STP 1.07 no FP8 speedup within noise little overlap CPX over SPX tok/s 0.78–0.96× – per GCD 1.76× launch-bound processes time (s) 3.01 – not comparable MI250X flushes FP16 TB/s 3.26 – 3.13×, 1.58× HBM bandwidth TFLOP/s 69.4 – 2.65×, 1.37× FP64 units

TABLE XI PARTITIONED AGAINST SHARED SERVING . Processes

partitioned

shared

TABLE XIII T RIAD BANDWIDTH AND ROC BLAS DGEMM RATE ON THE MI300A AND THE MI250X. T RIAD : a = b + s c (FP32), DGEMM: BEST SHAPE . ratio Triad, TB/s

MI300A, 6 processes: one per CPX partition vs all on SPX FP16 68.8 95.9 / 96.6 0.85 / 0.78 FP8 88.3 95.9 / 95.9 0.92 / 0.93 3 FP16 + 3 FP8 89.0 92.7 / 94.7 0.96 / 0.94 MI250X, 2 processes: one per GCD vs both on one GCD FP16 101.5 57.8 / 57.4 1.76 / 1.77

N

DGEMM, TFLOP/s

max (% peak) plateau peak run (% peak) sweep MI300A MI250X GCD vs MI300A MI250X pkg. vs MI300A

TABLE XII M EDIAN OF ROC HPL-M X P TIME TO SOLUTION AND RATES ON THE MI300A AND ONE MI250X GCD. C: ROC BLAS FP16 MEASURED PEAK ; RT : THE APPLICATION ’ S GEMM RATE ; IR SHARE : FRACTION OF TIME IN ITERATIVE REFINEMENT; nIR : REFINEMENT STEPS . t (s) R (TF/s) Φ=R/C Ψ=RT /C IR share nIR

MI300A (whole device) 30,720 0.192 100 61,440 0.732 211 112,640 3.01 316 138,240 5.12 344

Sec. V V VII VI VI VII VIII IX-A IX-A IX-A IX-B IX-C IX-C IX-D IX-D IX-D

0.178 0.373 0.559 0.608

0.672 0.682 0.742 0.753

0.260 0.164 0.096 0.076

3 3 3 3

MI250X, one GCD (FP16 LU flushed) 30,720 0.307 63.0 – 61,440 1.46 106 – 112,640 7.08 135 –

– – –

0.261 0.164 0.081

5 5 4

MI250X package (both GCDs at once, two solves; flushed) 30,720 0.308 126 – – 0.260 61,440 1.46 212 – – 0.164 112,640 7.09 269 – – 0.081

5 5 4

FP16 LU does its work, since multiplying the same operands in FP32 reproduces its residuals to within 0.43%. We therefore draw no generational ratio from this application, and the MI250X’s Φ = 1.05 and Ψ = 1.28, above its measured peaks, are the rate of a GEMM on zero operands. A generational comparison with an FP16 benchmark must first check that its operands are representable in FP16 on both GPUs. 1) Reference workloads: Table XIII lists a STREAM Triad, the same operation as BabelStream’s Triad [41], [42], and rocBLAS DGEMM. One GCD reaches 0.319 of the MI300A’s Triad bandwidth and 0.377 of its DGEMM rate, far more than its FP16 share in Table IV, so the MI250X is relatively stronger at FP64 and bandwidth than at FP16.

3.26 (61%) 1.04 (64%) 0.319 2.06 (63%) 0.633

3.20 0.98 0.307 1.99 0.621

69.4 (57%) 26.1 (55%) 0.377 50.7 (53%) 0.731

69.5 26.3 0.379 – –

X. D ISCUSSION Table X summarizes every workload on the MI300A: its result, what FP8 or 2:4 adds, its improvement over the MI250X, and the factor that limits it. Our measurements tell each audience what to change. Library developers. hipBLASLt’s heuristic costs up to 25% against the fastest kernel the library offers (Section VI), so kernel selection should rest on measurement. FP8 pays in an LLM only once its conversion is fused into the preceding kernel (Section IX-A), and hipSPARSELt keeps at most 0.554 of 2:4’s 2× because its kernels do not keep the matrix units supplied with data (Section VII). System administrators. Partitions isolate jobs that share an MI300A, but serve whole LLMs with less total throughput than SPX sharing (Section IX-C), while CU masks cannot confine a stream to fewer than 6 XCDs (Section VIII). The hardware-queue pool, not precision, sets fairness, and the default four queues trade fairness for throughput, so the pool size should follow whichever of the two matters. Users. FP8 is worth adopting for large, throughput-bound GEMMs, and dense should stay the default over 2:4. On the MI250X, which has neither feature, FP16 or BF16 through hipBLASLt is the choice, and moving that work to the whole MI300A gains 1.20–2.15× before FP8 is used. XI. C ONCLUSION Each loss we measured has a specific cause at a specific boundary between layers. We traced capability gain for FP8

and 2:4 sparsity on the MI300A against the MI250X, from instructions to applications and a shared device. FP8 keeps most of its speedup through the library, making an 81923 GEMM 3.95× faster than FP16 on the MI250X package. In an LLM, an unfused conversion to FP8 uses up that speedup (0.862× per layer), while fused serving keeps a gain at every load. 2:4 sparsity keeps at most about half of its speedup in hipSPARSELt, so dense wins. Under sharing, the hardwarequeue pool sets fairness and FP8’s GEMM capability gain rises above 1, but shared LLM serving gains nothing from FP8, and CPX partitions serve it with less throughput than SPX. Mixed-precision HPL yields no generational ratio, because the MI250X flushes its FP16 operands. R EFERENCES [1] P. Micikevicius, D. Stosic, N. Burgess, M. Cornea, P. Dubey, R. Grisenthwaite, S. Ha, A. Heinecke, P. Judd, J. Kamalu, N. Mellempudi, S. Oberman, M. Shoeybi, M. Siu, and H. Wu, “FP8 formats for deep learning,” 2022. [Online]. Available: https: //arxiv.org/abs/2209.05433 [2] Open Compute Project, “OCP 8-bit floating point specification (OFP8), revision 1.0,” Open Compute Project, Tech. Rep., Jun. 2023. [Online]. Available: https://github.com/opencomputeproject/FP8 [3] A. Mishra, J. A. Latorre, J. Pool, D. Stosic, D. Stosic, G. Venkatesh, C. Yu, and P. Micikevicius, “Accelerating sparse deep neural networks,” 2021. [Online]. Available: https://arxiv.org/abs/2104.08378 [4] J. Dongarra, P. Luszczek, and Y. Tsai, “HPL-MxP mixed-precision benchmark,” Benchmark website. [Online]. Available: https://hpl-mxp. org/ [5] A. Haidar, S. Tomov, J. Dongarra, and N. J. Higham, “Harnessing GPU tensor cores for fast FP16 arithmetic to speed up mixed-precision iterative refinement solvers,” in SC18: International Conference for High Performance Computing, Networking, Storage and Analysis. Dallas, TX, USA: IEEE, Nov. 2018, pp. 603–613. [Online]. Available: https://doi.org/10.1109/SC.2018.00050 [6] TOP500, “TOP500 list, november 2024,” https://top500.org/lists/top500/ 2024/11/, Nov. 2024. [7] ——, “TOP500 list, june 2026,” https://top500.org/lists/top500/2026/06/, Jun. 2026. [8] Introducing AMD CDNA 2 Architecture, Advanced Micro Devices, 2021, white paper. [Online]. Available: https://www.amd.com/content/dam/amd/en/documents/ instinct-business-docs/white-papers/amd-cdna2-white-paper.pdf [9] AMD Instinct MI200 Series Accelerator Data Sheet, Advanced Micro Devices, 2021. [Online]. Available: https://www.amd.com/content/dam/ amd/en/documents/instinct-tech-docs/instinct-mi200-datasheet.pdf [10] Introducing AMD CDNA 3 Architecture, Advanced Micro Devices, 2025, white paper. [Online]. Available: https://www.amd.com/content/dam/amd/en/documents/ instinct-tech-docs/white-papers/amd-cdna-3-white-paper.pdf [11] B. Hanindhito and B. Patel, “Characterizing performance, power, and energy of AMD CDNA3 GPU family,” in Proceedings of the International Conference for High Performance Computing, Networking, Storage and Analysis (SC ’25). ACM, 2025, pp. 905–934. [Online]. Available: https://doi.org/10.1145/3712285.3759768 [12] A. Tee, N. Curtis, N. Wolfe, and D. Wong, “The MALL is open: Exploring shared caches and latency in AMD CDNA 3 GPUs,” in Proceedings of the SC ’25 Workshops of the International Conference for High Performance Computing, Networking, Storage and Analysis. ACM, 2025, pp. 1110–1116. [Online]. Available: https://doi.org/10.1145/3731599.3767487 [13] AMD Instinct MI300A APU Data Sheet, Advanced Micro Devices, 2025. [Online]. Available: https: //www.amd.com/content/dam/amd/en/documents/instinct-tech-docs/ data-sheets/amd-instinct-mi300a-data-sheet.pdf [14] “AMD Instinct MI300” Instruction Set Architecture Reference Guide, Advanced Micro Devices, Aug. 2025, cDNA3 ISA; revision of 5 August 2025. [Online]. Available: https://www.amd.com/content/ dam/amd/en/documents/instinct-tech-docs/instruction-set-architectures/ amd-instinct-mi300-cdna3-instruction-set-architecture.pdf

[15] A. Abouelmagd, D. Boehme, S. Brink, J. Burmark, M. McKinsey, A. Skjellum, and O. Pearce, “GPU partitioning, power, and performance of the AMD MI300A,” in Proceedings of the Supercomputing Asia and International Conference on High Performance Computing in Asia Pacific Region (SCA/HPCAsia ’26). ACM, 2026, pp. 215–227. [Online]. Available: https://doi.org/10.1145/3773656.3773680 [16] HSA Runtime Programmer’s Reference Manual, HSA Foundation, 2021. [Online]. Available: https://hsafoundation.com/standards/ [17] G. Schieffer, D. Araújo De Medeiros, J. Faj, A. Marathe, and I. Peng, “On the rise of AMD matrix cores: Performance, power efficiency, and programmability,” in 2024 IEEE International Symposium on Performance Analysis of Systems and Software (ISPASS). IEEE, 2024, pp. 132–143. [Online]. Available: https: //doi.org/10.1109/ISPASS61541.2024.00022 [18] A. Jarmusch and S. Chandrasekaran, “Microbenchmarking NVIDIA’s Blackwell architecture: An in-depth architectural analysis,” in 2026 IEEE International Parallel and Distributed Processing Symposium (IPDPS). IEEE, 2026, pp. 1026–1036. [Online]. Available: https: //doi.org/10.1109/IPDPS65963.2026.00087 [19] C. Ambati and T. Diep, “AMD MI300X GPU performance analysis,” 2025. [Online]. Available: https://arxiv.org/abs/2510.27583 [20] W. Luo, R. Fan, Z. Li, D. Du, Q. Wang, and X. Chu, “Benchmarking and dissecting the Nvidia Hopper GPU architecture,” in 2024 IEEE International Parallel and Distributed Processing Symposium (IPDPS). IEEE, 2024, pp. 656–667. [Online]. Available: https://doi.org/10.1109/IPDPS57955.2024.00064 [21] H. Wong, M.-M. Papadopoulou, M. Sadooghi-Alvandi, and A. Moshovos, “Demystifying GPU microarchitecture through microbenchmarking,” in 2010 IEEE International Symposium on Performance Analysis of Systems & Software (ISPASS), 2010, pp. 235– 246. [Online]. Available: https://doi.org/10.1109/ISPASS.2010.5452013 [22] X. Mei and X. Chu, “Dissecting GPU memory hierarchy through microbenchmarking,” IEEE Transactions on Parallel and Distributed Systems, vol. 28, no. 1, pp. 72–86, 2017. [Online]. Available: https://doi.org/10.1109/TPDS.2016.2549523 [23] J. Wahlgren, G. Schieffer, R. Shi, E. A. León, R. Pearce, M. Gokhale, and I. Peng, “Dissecting CPU-GPU unified physical memory on AMD MI300A APUs,” in 2025 IEEE International Symposium on Workload Characterization (IISWC). IEEE, Oct. 2025, pp. 368–380. [Online]. Available: https://doi.org/10.1109/IISWC66894.2025.00038 [24] G. K. R. Kuncham, S. Zhang, and D. K. Panda, “One memory-many paths: Early experiences with allocation and data copy strategies on MI300A,” in 2026 IEEE International Parallel and Distributed Processing Symposium (IPDPS). IEEE, 2026, pp. 1012–1025. [Online]. Available: https://doi.org/10.1109/IPDPS65963.2026.00086 [25] O. Antepara, L. Oliker, and S. Williams, “Roofline analysis of tightly-coupled CPU-GPU superchips: A study on MI300A and GH200,” in Proceedings of the SC ’25 Workshops of the International Conference for High Performance Computing, Networking, Storage and Analysis. ACM, 2025, pp. 1205–1216. [Online]. Available: https://doi.org/10.1145/3731599.3767497 [26] G. Schieffer, J. Wahlgren, R. Shi, E. A. León, R. Pearce, M. Gokhale, and I. Peng, “Inter-APU communication on AMD MI300A systems via Infinity Fabric: A deep dive,” in Proceedings of the International Symposium on Memory Systems (MEMSYS ’25). ACM, 2025, pp. 238–250. [Online]. Available: https://doi.org/10.1145/3767110.3767130 [27] N. Otterness and J. H. Anderson, “AMD GPUs as an alternative to NVIDIA for supporting real-time workloads,” in 32nd Euromicro Conference on Real-Time Systems (ECRTS 2020), ser. Leibniz International Proceedings in Informatics (LIPIcs), vol. 165. Schloss Dagstuhl–Leibniz-Zentrum für Informatik, 2020, pp. 10:1–10:23. [Online]. Available: https://doi.org/10.4230/LIPIcs.ECRTS.2020.10 [28] ——, “Exploring AMD GPU scheduling details by experimenting with “worst practices”,” in Proceedings of the 29th International Conference on Real-Time Networks and Systems (RTNS ’21). ACM, 2021, pp. 24–34. [Online]. Available: https://doi.org/10.1145/3453417.3453432 [29] R. Quach, Y. Wang, A. Jahanshahi, D. Wong, and H. Kim, “ECLIP: Energy-efficient and practical co-location of ML inference on spatially partitioned GPUs,” in 2025 IEEE/ACM International Symposium on Low Power Electronics and Design (ISLPED). IEEE, 2025, pp. 1–7. [Online]. Available: https://doi.org/10.1109/ISLPED65674.2025. 11261793 [30] A. Masood, P. Gaur, and N. Jayasena, “RAPID-Serve: Resource-

efficient and accelerated P/D intra-GPU disaggregation,” 2026. [Online]. Available: https://arxiv.org/abs/2601.11822 [31] W. Wen, C. Wu, Y. Wang, Y. Chen, and H. Li, “Learning structured sparsity in deep neural networks,” in Advances in Neural Information Processing Systems 29 (NIPS 2016), 2016, pp. 2074–2082. [Online]. Available: https://arxiv.org/abs/1608.03665 [32] Y. Hu, K. Zhao, W. Huang, J. Chen, and J. Zhu, “Accelerating transformer pre-training with 2:4 sparsity,” in Proceedings of the 41st International Conference on Machine Learning (ICML), ser. Proceedings of Machine Learning Research, vol. 235, 2024, pp. 19 531–19 543. [Online]. Available: https://arxiv.org/abs/2404.01847 [33] Y. Hu, J. Zhu, and J. Chen, “S-STE: Continuous pruning function for efficient 2:4 sparse pre-training,” in Advances in Neural Information Processing Systems 37 (NeurIPS 2024), 2024. [Online]. Available: https://arxiv.org/abs/2409.09099 [34] J. Cai, D. Haziza, and S. Rao, “Accelerating neural network training with semi-structured (2:4) sparsity,” PyTorch Blog, Jun. 2024. [Online]. Available: https://pytorch.org/blog/accelerating-neural-network-training/ [35] D. Dickmann, P. Offenhäuser, R. Saxena, G. Markomanolis, A. Rigazzi, P. Keller, K. Kayabay, and D. Hoppe, “Evaluating AMD Instinct MI300A APU: Performance insights on LLM training via knowledge distillation,” in Proceedings of the Cray User Group (CUG ’25). ACM, 2025, pp. 115–126. [Online]. Available: https://doi.org/10.1145/ 3757348.3757361 [36] W. Kwon, Z. Li, S. Zhuang, Y. Sheng, L. Zheng, C. H. Yu, J. E. Gonzalez, H. Zhang, and I. Stoica, “Efficient memory management for large language model serving with PagedAttention,” in Proceedings of the 29th Symposium on Operating Systems Principles (SOSP), 2023, pp. 611–626. [37] Advanced Micro Devices, “rocHPL-MxP: HPL-MxP benchmark for ROCm,” GitHub repository. [Online]. Available: https://github.com/ ROCm/rocHPL-MxP [38] hipSPARSELt Documentation, Advanced Micro Devices, 2026, accessed September 2026. [Online]. Available: https://rocm.docs.amd.com/ projects/hipSPARSELt/en/latest/ [39] R. Jain, D.-M. Chiu, and W. Hawe, “A quantitative measure of fairness and discrimination for resource allocation in shared computer systems,” Digital Equipment Corporation, Tech. Rep. DEC-TR-301, Sep. 1984. [Online]. Available: https://arxiv.org/abs/cs/9809099 [40] rocBLAS Documentation, Advanced Micro Devices, 2026, accessed September 2026. [Online]. Available: https://rocm.docs.amd.com/ projects/rocBLAS/en/latest/ [41] T. Deakin, J. Price, M. Martineau, and S. McIntosh-Smith, “GPUSTREAM v2.0: Benchmarking the achievable memory bandwidth of many-core processors across diverse parallel programming models,” in High Performance Computing: ISC High Performance 2016 International Workshops, Revised Selected Papers, ser. Lecture Notes in Computer Science, M. Taufer, B. Mohr, and J. M. Kunkel, Eds. Cham: Springer International Publishing, 2016, pp. 489–507. [Online]. Available: https://doi.org/10.1007/978-3-319-46079-6 34 [42] ——, “Evaluating attainable memory bandwidth of parallel programming models via BabelStream,” International Journal of Computational Science and Engineering, vol. 17, no. 3, 2018. [Online]. Available: https://doi.org/10.1504/IJCSE.2018.095847

Related documents

• How to grade the accuracy of the BLAS#1000265
• CRISP: Co-designing Pruning and Scheduling for Efficient MoE Inference on Edge Servers#1473611
• CRISP: Co-designing Pruning and Scheduling for Efficient MoE Inference on Edge Servers#1469045
• LRMoE: Demand-Side Load Re-Calibration for Efficient MoE Training on Memory-Constrained Nodes#1488512
• LRMoE: Demand-Side Load Re-Calibration for Efficient MoE Training on Memory-Constrained Nodes#1482981