Conceptio › Archive › arXiv CS

arXiv CSopen access

Microbenchmark-Driven Analytical Performance Modeling Across Modern GPU Architectures

2026 · arxiv_cs

arXiv CS · Papers · License: Open Access · 2026

Open Source ↗Direct PDF ↓

cloud distributed-computing parallel-computing

distributed computing, parallel computing, cloud

arXiv:2605.04178v1 [cs.DC] 5 May 2026

Microbenchmark-Driven Analytical Performance Modeling Across Modern GPU Architectures Aaron Jarmusch

Sunita Chandrasekaran

dept. Computer & Information Sciences University of Delaware Newark, Delaware, USA

Abstract—Rapidly evolving GPU architectures featuring complex memory hierarchies, matrix units, and varied precision formats continue to widen the gap between theoretical peaks and achievable performance. We design and develop analytical performance models for NVIDIA Blackwell (B200) and AMD CDNA3 (MI300A) grounded in systematic microbenchmark characterization. For Blackwell, the model captures Tensor Memory (TMEM), asynchronous bulk copy (TMA), and 5th-generation tensor cores; for CDNA3, the model captures Infinity Cache hierarchy, VGPR constraints, and occupancy. Validation yields 1.31% MAE on B200 (21 kernels) and ∼0.09% on MI300A (27 kernels), while naive roofline baselines exceed 95% error on the same kernels. We further validate the models using Rodinia 3.1 and SPEChpc 2021 Tiny.The models are updated with HBM bandwidth, capacity, and cache parameters and applied to H200 (Hopper) and MI250X (CDNA2), indicating no major restructuring of the models are needed. All models and benchmarks will be released as open-source upon acceptance. Index Terms—GPU performance modeling, analytical model, microbenchmark, Blackwell, CDNA3, MI300A, roofline, HPC

I. I NTRODUCTION Modern HPC and AI systems increasingly rely on GPU accelerators that expose distinct execution primitives across vendors and generations (For example, Tensor Memory (TMEM) and asynchronous bulk copy (TMA) on NVIDIA Blackwell, wavefront scheduling and Infinity Cache on AMD CDNA), and whose sustained performance diverges significantly from datasheet peaks [1]–[3]. Every 18–24 months a new architecture is introduced with new memory subsystems, matrix execution units, and precision formats. Understanding the gap between theoretical peak and achievable efficiency requires architecture-aware models that capture memory hierarchy behavior, matrix unit utilization, and occupancy constraints, not just bandwidth-limited roofline bounds. This paper presents systematic microbenchmark-driven stage-centric analytical models for two current-generation GPU accelerators: NVIDIA Blackwell (B200) [4] and AMD CDNA3 (MI300A) [5]. Stage-centric implies the different stages in a GPU compilation. Our models are build on prior microbenchmark characterizations of both architectures [6], [7]. We characterize each architecture through targeted lowlevel benchmarks, derive model parameters directly from measurements, and validate against both a microbenchmark suite and two HPC application benchmarks (Rodinia 3.1 [8] and SPEChpc 2021 Tiny [9]). The models are interpretable,

parameterized by measured hardware values, and accurate to within 1–5% MAE, compared to over 95% error for naive roofline. We then apply the same methodology to the prior generation of each vendor (NVIDIA Hopper H200 and AMD CDNA2 MI250X), demonstrating that both ports required only hardware parameter updates with no model re-derivation. We make the following novel contributions: • Construct and validate a stage-centric analytical model for NVIDIA Blackwell (B200) capturing TMA, TMEM, 5th-generation tensor cores, and the 2-SM cooperative execution model; to our knowledge the first validated execution-time model for this architecture • Construct and validate stage-centric analytical model for AMD CDNA3 (MI300A) accounting for Infinity Cache hierarchy, VGPR register pressure, and occupancy-driven tile selection • Validate both models against Rodinia 3.1 and SPEChpc 2021 Tiny on B200 and MI300A, with naive roofline baselines showing >95% error • By updating parameters of our analytical models, they can be applied to H200 and MI250X. Paper organization. Section II positions our work against prior analytical and ML-based models. Section III summarizes Blackwell and CDNA architectures and their implications for modeling. Section IV presents the stage-centric model for Blackwell and the wavefront-centric model for AMD CDNA. Section V describes the microbenchmark validation setup, results, and cross-platform comparison with naive roofline baselines. Section VI discusses systematic bias, limitations, architectural insights, and portability. Section VII presents the conclusion. II. R ELATED W ORK Related work spans roofline modeling, analytical and simulation-based GPU performance prediction, and microarchitecture characterization. A. Roofline and Analytical Models The roofline model [10] relates arithmetic intensity to attainable performance, identifying memory-bound vs. computebound regimes. Cache-aware extensions [11], [12] and AMD instruction-roofline variants [13] improve diagnosis, but all

B. Simulation, Characterization, and ML Approaches Accel-Sim [18] offers cycle-accurate simulation but depends on architecture definition files often unavailable for the newest accelerators. Analytical models trade accuracy on irregular control flow for interpretability and negligible evaluation cost. Key characterization work: Wong et al. [19] established GPU microbenchmarking methodology; Jia et al. [20] dissected Volta; Luo et al. [21] characterized Hopper TMA/wgmma/DSM (basis for our H200 calibration). On AMD, Wahlgren et al. [22] studied MI300A unified memory and Schieffer et al. [23] characterized matrix-core performance. Jarmusch et al. [6] provided the first Blackwell (B200) microbenchmark characterization; Jarmusch et al. [7] characterized MI300A FP8 matrix cores and ACE concurrency. Our analytical models build directly on these measurements. Chowdhury et al. [24] modeled tensor core units; AMALI [25] targets LLM inference; Fasi et al. [26] analyze tensor-core numerics. None deliver full-kernel analytical models with reported MAE across vendors. C. Positioning Our work is, to our knowledge, the first to (a) provide validated analytical execution-time models for Blackwell (B200) and CDNA3 (MI300A), (b) update parameters of our analytical models and apply to H200 and MI250X, and (c) report cross-vendor MAE under a shared protocol. Prior models require full re-derivation or simulator updates when new hardware ships; we close this gap by tying every coefficient to a microbenchmark. Table I summarizes. III. GPU A RCHITECTURE OVERVIEW We focus on the NVIDIA Blackwell B200 [4] and AMD CDNA 3 MI300A [5]; two state-of-the-art architectures rep-

Chip Die #1

Chip Die #2

SM SM SM SM

L0 Instruction Cache

Warp Scheduler (32 thread/cclk) L0 (32 Instruction Cache Dispatch Unit thread/cclk)

Graphics Processor Cluster (GPC)

MIG Control

Graphics Processor L2 Cache Cluster (GPC)

Warp Scheduler (32 thread/cclk)

Graphics Processor Cluster (GPC)

Warp Scheduler (32 thread/cclk)

Tensor Memory () Register File(TMEM) (16,384 xthread/cclk) 32-bit) L0 (32 Instruction Cache Dispatch Unit Warp Scheduler (32 thread/cclk)

Tensor Memory (TMEM) () 32-bit) Register File (16,384 CUDA Dispatch Unit (32xthread/cclk) 5TH-Gen Tensor Cores Cores Tensor Memory (TMEM) Register File (16,384 () x 32-bit) CUDA 5TH-Gen LD/ST

NV-HBI

Graphics Processor L2 Cache Cluster (GPC)

NV-HBI

remain single-axis intensity plots that cannot capture per-stage pipeline overlap, dedicated matrix paths, or TMEM residency. Why roofline gives >95% error on modern kernels. The naive bound Troofline = max(FLOPs/Ppeak , bytes/BHBM ) fails for three compounding reasons: (1) datasheet peaks overstate achievable throughput (B200 sustained tensor-core throughput is 1,100–1,400 TFLOPS vs. 2,250 datasheet; sustained HBM is 6.8–7.1 TB/s vs. 8.0), making Troofline too small by 1.5–2×; (2) roofline ignores serialized pipeline stages (on Blackwell, TMA→TMEM→Tensor Core stages add latency that max() cannot represent); (3) roofline uses a single bandwidth figure, missing MI300A’s 256 MB Infinity Cache (17.2 TB/s vs. 5.3 TB/s HBM). These compound multiplicatively to >90% error (Table VI). Hong and Kim [14] established the MWP/CWP framework for bounding parallelism on early NVIDIA GPUs; we retain that occupancy intuition but add Blackwell-specific TMEM/TMA/sync terms. GCoM [15] provides a detailed Ampere-era core model; we target newer architectures. GPUMech [16] and MDM [17] address interval analysis and memory divergence respectively, neither targeting tensor cores or TMEM.

Tensor Cores Cores Tensor Memory (TMEM) () SFU CUDA

Cores

LD/ST SFU L1 Instruction Cache CUDA

Cores

5TH-Gen Tensor Cores

TMA LD/ST SFU L1 Instruction Cache

Graphics Processor Cluster (GPC)

5TH-Gen Tensor Cores

TMA LD/ST SFU L1L1 Instruction Cache 256 KB Data Cache/Shared Memory TMA L1L1 Instruction Cache 256 KB Data Cache/Shared Memory TMA 256 KB L1 Data Cache/Shared Memory 256 KB L1 Data Cache/Shared Memory

Fig. 1: NVIDIA Blackwell architecture: dual-die, TMEM, and SM structure. CCD

XCD

L2 x38 CUs

CCD CCD

Infinity Fabric Infinity Cache

Infinity Cache

HBM

Fig. 2: MI300A APU: GPU/CPU chiplets and unified memory. resenting the leading edge of heterogeneous integration and matrix acceleration. NVIDIA Blackwell (B200). Dual-die design (208B transistors) with NV-HBI (10 TB/s inter-die), unified 192 GB HBM3e and cache coherence [4]. Fifth-gen Tensor Cores support FP4/FP6 with up to 9,000 TFLOPS FP4; Transformer Engine improves low-precision stability. Tensor Memory (TMEM) provides 256 KB/SM for tensor ops, reducing shared-memory contention. A decompression engine supports LZ4/Snappy/Deflate in the data path. The Tensor Memory Accelerator (TMA) handles hardware-managed async transfers; tcgen05.mma integrates TMEM and weight-stationary dataflows (Figure 1). AMD MI300A. First heterogeneous APU: six GPU chiplets (GCDs) and three Zen 4 CPU chiplets with 24 cores and Unified Physical Memory (UPM)—true hardwarecoherent CPU–GPU sharing without explicit copies [5]. Matrix Cores support FP8 (1,307 TFLOPS) and strong FP64 (61.3 TFLOPS); 256 MB L2 (Infinity Cache) reduces latency for large working sets. Execution is wavefront-based (64 threads); 304 CUs (38 per XCD × 8 XCDs) with TF32 and 2:4 sparsity. Cross-XCD access incurs 50–100 ns NUMA-like penalties (Figure 2). Modeling implications. Blackwell’s explicit stages (TMA → TMEM → Tensor Core → Sync) allow stage-centric modeling with measurable latencies and bandwidths; roofline’s single max() cannot represent their serialization. MI300A’s overlap is implicit and occupancy-driven; accumulators live in VGPRs (vs. TMEM), creating tile-size vs. occupancy tradeoffs that a bandwidth roofline cannot capture. These structural differences are why a single roofline formulation gives >95% error on both platforms (Table VI), and why the two architectures require distinct model frameworks rather than

TABLE I: Related work vs. this paper: coverage of modern features, cross-vendor validation, and porting methodology. Work Hong et al. [14] GCoM [15] Chowdhury et al. [24] Wahlgren/Schieffer [22], [23] This work

Arch.

TMEM/TMA/Decomp.

NVIDIA NVIDIA Tensor Cores AMD MI300A Blackwell, CDNA 3

– – TC only – Yes

Time SMEM

TMA Load 0

TMEM

Compute

TMA Load 1

TMA Load 2

TMA Load 3

tcgen05.cp

tcgen05.mma

HBM

Total Execution Time (Critical Path)

Fig. 3: Per-CTA execution: pipeline stages, overlap, and critical path. parameter substitution in a shared roofline. They also illustrate the principle of extensibility. If a future GPU introduces a distinct accumulation mechanism—for instance, dedicated register files for FP4 on a future CDNA4 or NV-HBI-aware TMEM on Rubin integration is a matter of identifying the most similar framework and adding the new term without the need to rebuild the model from the ground up. Table II makes this concrete: every parameter is either measured by a microbenchmark (bandwidth, throughput, max resident warps) or taken from the vendor datasheet (SM/CU count, cache sizes, capacity). Swapping in values for a new GPU updates the model without changing any formula. IV. A NALYTICAL M ODEL We adopt the Hong–Kim framework [14]: execution time is the maximum of compute and memory time plus overhead: Texec = max(Tcompute , Tmemory ) + Toverhead

(1)

Toverhead covers barriers and kernel launch. Blackwell [4] uses explicit stage-centric pipelines; MI300A [5] uses implicit wavefront-centric scheduling. Figure 3 summarizes the perCTA pipeline for Blackwell.

1) Tensor Core Compute and TMEM: TMEM (256 KB/SM) holds accumulators; access has measurable bandwidth. Per K-step tile iteration: Daccum Daccum + Lmma + (2) BWTMEM read BWTMEM write

Exceeding 256 KB forces spill and hurts efficiency. Per-CTA compute time: Tcompute =

2bM bN bK + TTMEM + TTMEM mgmt RTSM C × Smode

Validated MAE

No No No No Yes

– Yes – Partial 1.3% (B200), ∼0.09% (MI300A)

RTSM C is tensor core throughput per SM; Smode accounts for 2-SM cooperation. Two TMEM strategies: (A) accumulators only in TMEM, inputs in SMEM—lower TMEM traffic; (B) A and accumulators in TMEM via tcgen08.cp—reduces SMEM pressure but increases TMEM bandwidth. The model covers both via TTMEM and measured bandwidths. 2) Memory and TMA: TMA multicast: P participants share a tile; bytes per CTA: bytesperCT A = bytes(T )/P . TMA time per CTA: bytes(T ) (4) Ttma = LT M A + P × BT M A LT M A , BT M A from microbenchmarks; L2 hit rates strongly affect BT M A . For different PA , PB on A/B, sum or take max depending on overlap. 3) Decompression and Synchronization: = Decompression (link vs. engine): TDE load max(Dcompressed /BWlink , Dcompressed /RDE ). With compression ratio CR and efficiency ηDE : TDE load =

Duncompressed CR × BWlink × ηDE

(5)

Sub-byte unpacking: Tdecomp = bytescomp /Rdecomp + Ldecomp setup ; Rdecomp and overlap via α. We use α ∈ [0.85, 0.95] from pipeline depth (double- to triple-buffering); sensitivity is discussed in Section V. Sync per K-step: Tsync = Nbar × Lmbar ; Nbar typically 1–2, Lmbar from microbenchmarks. 4) CTA Pairs (2-SM UMMA): CTA pairs on adjacent SMs share B via distributed shared memory (DSMEM). Traffic: D2−CT A = 2MA + MB (vs. 2(MA + MB )), giving up to ∼1.33× traffic reduction for square tiles. Memory time Tmemory 2−CT A = (2MA + MB )/BWshared . Sync cost Ktiles × Lcommit . Compute with 2-SM: 2bM bN bK + TT M EM + TT M EM mgmt RTSM C × S2SM (6) S2SM is measured speedup (Section V). 5) Overlap and Critical Path: Overlap factor α ∈ [0, 1] (fraction of I/O hidden behind compute): Tcompute 2SM =

A. NVIDIA Blackwell Model

TTMEM per tile =

Cross-vendor

(3)

ef f Tio = (1 − α)(Ttma + Tdecomp ) + Tsync

(7)

Writeback: Tstore = bytes(Ctile )/Bgmem + Lstore setup or TMA store LT M A store + bytes(Ctile )/BT M A ; often overlapped in persistent kernels. TMEM management amortized: TTamortized M EM mgmt = (Lalloc + Ldealloc )/Ktiles .

TABLE II: Architecture parameters for B200 and MI300A. Each value is from a microbenchmark or vendor datasheet (Source column). H200 and MI250X use the same model with their own values. FP64 roofline for SPEChpc uses 30.4 TFLOPS on MI300A. Parameter

B200

MI300A

Source

SMs / CUs Warp / wavefront size Max resident warps/wavefronts HBM peak BW (TB/s) HBM capacity (GB) L2 / LLC (MB) TMEM / LDS (KB per SM/CU) Tensor / MFMA peak (TFLOPS)

176 32 64 8.0 192 64 256 2,250 (FP16), 4,500 (FP8)

304 64 32 5.3 128 256 64 1,307 (FP8), 61.3 (FP64)

Datasheet Datasheet Microbench / docs Bandwidth microbench / datasheet Datasheet Datasheet Datasheet Throughput microbench / datasheet

Steady-state pipelined step: Tstep pipelined max(Ttma , Tdecomp , Tcompute , Tsync ) + ϵ. Per-step time: ef f Tstep = max(Tcompute , Tio ) + Tsync + Omisc

(8)

Omisc includes TMEM mgmt and pipeline bubbles. Total kernel time: Ktiles × Tstep plus launch and writeback. 6) Concurrent streams and multi-GPU: For Nc concurrent streams, we add (Nc −1)τc to single-stream time; for Nd GPUs, (Nd −1)τg . Both are fit from microbenchmarks, mirroring the MI300A interference terms. B. AMD MI300A Model Overlap is occupancy-driven; memory goes through L1→L2→LLC→HBM (latencies measured via pointerchasing [27]); accumulators in vector general-purpose registers (VGPRs). Overlap factor: active (Nwf − 1) × Tcompute ηoverlap = min 1.0, (9) Tmemory Effective memory time (hit rates hL1 , hL2 , hLLC ):

Effective bandwidth: BWeffective = hLLC · BWLLC + (1−hLLC ) · BWHBM . The LLC hit rate hLLC (W ) is a piecewise function of working set W (MB), summarized in Table III. TABLE III: MI300A Infinity Cache hit rate model hLLC (W ). hLLC (W )

Regime

W < 205 MB

1.0 α 1 − W −205 51 256 β

Transition zone

205 ≤ W ≤ 256 MB W > 256 MB

Fully cache-resident Streaming / spill to HBM

Here α and β capture access pattern and streaming behavior (calibrated from microbenchmarks). Matrix Fused MultiplyAdd (MFMA) compute: MFMA Tcompute =

NMFMA inst NCU × ThroughputMFMA × Utilization

MI300A Tstep =

active Nwf Per-step (overlap

eff MFMA Tmemory + Tcompute 1 + ηoverlap

MI300A MI300A Tkernel = Tlaunch + Ktiles × Tstep

+Twriteback + Tcoherence + Tcross XCD

(11)

= in

(12)

(13)

Coherence and cross-XCD account for unified memory and NUMA. For MI300A validation we derive FLOPs and bytes for each case directly from the real problem sizes (M , N , K, vector length N , density, etc.), using synthetic entries only as templates for fixed micro-architectural counts. For the occupancy/tile study (8×8 vs. 16×16 tiles), we use a pipeline/occupancy model for MI300A: per-CTA step time Tstep cta = max(flops per cta/peak, bytes per cta/BWeff ) with hLLC (W ) from working set. Total kernel time includes a scheduling term: Nctas · Tstep cta NCU · Weff +Twriteback + Tcoherence + Tcross XCD ,

occ Tkernel = Tlaunch + τcta · Nctas +

eff Tmemory = Nloads × hL1 LL1 + (1 − hL1 )hL2 LL2 +(1 − hL1 )(1 − hL2 )hLLC LLLC + (1 − htotal )LHBM (10)

Working set

VGPR-limited occupancy: min(32, ⌊65536/VGPRper wf ⌋). denominator):

(14)

where Nctas is the grid CTA count, Weff is effective wavefronts per CU, and τcta is overhead per CTA (tunable from validation). This formulation yields the correct ordering (16×16 faster than 8×8); Weff and τcta can be tuned to match measured runtimes. The model extends MI300A with optional memory warp parallelism (MWP) and compute warp parallelism (CWP) eff limits [14]: effective wavefronts for overlap are Nwf = active min(Nwf , MWP, CWP) when MWP and CWP are set (perCU limits from microbenchmarks or tuning); ηoverlap in Eq. 9 eff then uses Nwf . The validation MAE reported in Section V uses the base model (MWP/CWP not set). When multiple kernels run concurrently (e.g., on different HIP streams), we add an interference term to Eq. 13: multi MI300A Tkernel = Tkernel + (Nconcurrent − 1) · τinterf , where Nconcurrent is the number of concurrent kernels and τinterf is the measured overhead per additional concurrent kernel (seconds). Singlekernel (Nconcurrent = 1) is unchanged. Validation results (Section V) use the tuned τinterf = 50 µs.

For runs using multiple MI300A devices, we add an optional multi-GPU MI300A multi-GPU term: Tkernel = Tkernel +(Ndevices −1)·τinterf gpu , where Ndevices is the number of GPUs and τinterf gpu is the measured overhead per additional GPU (seconds). Single-GPU (Ndevices = 1) is unchanged. Copy and sync between devices are not modeled separately and can be absorbed into τinterf gpu when tuned from measurements. The model supports adaptive tile selection (evaluate candidate tiles via Eq. 13 and return the minimum-time tile) and kernel fusion (combined FLOPs/bytes plus optional overhead τfusion ). These predict relative tile costs and fused-kernel runtime, not the compiler’s internal tile choice. Apply models to H200 and MI250X. The H200 uses the same Hopper roofline structure with HBM bandwidth of 4.8 TB/s and capacity of 141 GB; no new model terms are required. The MI250X uses the same CDNA framework as MI300A with its own peak FP64 (383 TFLOPS), bandwidth (3.2 TB/s), and cache hierarchy (128 MB LLC, 220 CUs); the occupancy/tile cases use calibrated scaling analogous to MI300A. Section V reports results for both targets. C. Unified Summary Table IV summarizes how to compute execution time and which parameters drive each term; use it by characterizing the workload first (arithmetic intensity, working set, access pattern), then instantiating the appropriate row per architecture. TC , Tioeff ) + Blackwell: TBlackwell = Tlaunch + Ktiles × max(Tcompute eff Tsync + Twriteback . MI300A: Eq. 13 with Tmemory , ηoverlap from occupancy and hLLC (W ). D. Model Workflow and Calibration To apply the model: (1) characterize the workload (arithmetic intensity, working set W , tile dimensions, class); (2) select parameters from Table VII or Table II; (3) apply the appropriate formula (Blackwell: TBlackwell = Tlaunch + Ktiles × TC max(Tcompute , Tioeff ) + Tsync + Twriteback ; MI300A: Eq. 13). Optionally provide precision (fp16/fp8/fp64) for tensor efficiency and tile sizes for tile-aware paths. Example: GEMM with M =N =K=16384 on B200, tile 128×128×32, predicts 4.17 ms (measured: 4.10 ms, 1.8% error). Calibration. First-principles parameters (bandwidths, Tlaunch , barrier latencies) come from microbenchmarks. Optional per-case multipliers may align predictions with profiler kernel-sum times; such factors must be disclosed. We recommend train/holdout splits when calibration is used. E. Host–Device Transfer and Synchronization When validation or deployment concerns host–device traffic or explicit host synchronization (cudaDeviceSynchronize / hipDeviceSynchronize), the implementation extends the generic segment schema with transfer and synchronization phases. For each transfer episode moving S bytes between host and device, Tmemcpy =

S + τmemcpy , dir Beff

(15)

dir where Beff is effective bandwidth (bytes/s) for H2D or D2H (defaults are conservative; measured with platform microbenchmarks or overridden at run time). The fixed term τmemcpy amortizes API launch overhead. Each counted synchronization point contributes Thost sync = τsync . Segment times multiply by nexec in the workload segment file like other phases. Overlap between copy and kernel execution is not modeled in this version; the sum is conservative when compared to wall-clock overlap.

F. Generic roofline path: scaling, working set, and launches When a segment does not map to a full Blackwell stage model or a validated GEMM/tile case, the implementation uses a generic roofline with separate calibrated scales for memory-bound, compute-bound, balanced, and stencil classes, optional precision-specific tensor efficiency multipliers, and a working-set–aware global memory bandwidth Beff (W ) = Bsustained + (Bpeak − Bsustained ) exp(−W/w0 ), (16) with w0 a tunable working-set scale (set ≤ 0 to disable the blend). This captures the idea that smaller resident working sets can see higher effective bandwidth than streaming from a footprint that saturates HBM. For multi-kernel segments, each workload row may specify multiple kernels; the model adds extra launch time beyond the first kernel’s modeled time, using measured launch latency from the same configuration as the generic path. G. Assumptions and Extensions The model assumes regular memory access, known or estimated cache hit rates, single-kernel single-GPU execution (by default), and steady-state frequency. Required inputs: for Blackwell, bM , bN , bK , Ktiles , bytes per CTA, TMA participants P , α; for MI300A, tile dimensions, Ktiles , bytes, hL1 /hL2 /hLLC (W ), occupancy. Optional extensions include MWP/CWP limits, LDS bank conflicts, multi-kernel/multiGPU interference, adaptive tile selection, and kernel fusion. Not yet modeled: CTA queuing delays and multi-node scaling. V. M ODEL VALIDATION AND ACCURACY This section presents the microbenchmark validation methodology, per-platform results, and two application case studies (Rodinia 3.1 and SPEChpc 2021 Tiny). A. Microbenchmark Design Model parameters are derived from a custom microbenchmark suite. For Blackwell: we measure (i) TMEM read/write bandwidth via tile copy between TMEM and SMEM; (ii) TMA copy latency and effective bandwidth as a function of tile size and L2 residency; (iii) mbarrier wait and commit latency (Lmbar , Lcommit ); (iv) tensor core instruction latency (tcgen05.mma, tcgen08.cp) and peak throughput by precision. For MI300A: we measure (i) L1/L2/Infinity Cache and HBM bandwidth and latency (latency vs. outstanding warps, bandwidth ceilings); (ii) MFMA throughput and utilization vs. tile size; (iii) hLLC (W ) via sweep over working set size W .

TABLE IV: Model summary: Blackwell vs. MI300A Component

Blackwell B200

AMD MI300A

Execution Compute Memory Overlap Sync Constraints

max(Tcompute , Tioeff ) + Tsync Tensor core + TMEM (Eq. 3) TMA (Eq. 4), α overlap α pipeline depth Explicit barriers Lmbar TMEM 256 KB/SM

eff + Tcompute )/(1 + ηoverlap ) (Tmemory MFMA (Eq. 11), Util 0.4–0.7 Cache hierarchy (Eq. 10), hLLC (W ) ηoverlap occupancy (Eq. 9) Implicit; coherence 100–200 ns VGPR → occupancy

MAE (typ.)

5.4–8.4% (B200)

∼0.09% (MI300A n=27; calibrated)

TABLE V: Measurement protocol metadata (recorded on the MI300A Rodinia validation host). Item

Recorded value / policy

NVIDIA GPU AMD GPU Driver / firmware CUDA / toolkit Power / clocks Microbenchmark list Validation apps

SKU, board ID if available, stepping SKU, socket/APU vs discrete if applicable NVIDIA driver build, ROCm version Version used to build microbenchmarks and apps Persistence mode, locked clocks, or default policy Same binaries and flags across all validation runs NVIDIA: Nsight Systems; AMD: ROCm rocprof

a) Measurement protocol and platform lock-in.: Validation and microbenchmarks use the same GPU stepping, driver stack, and power/clock policy unless a sensitivity study explicitly varies them. Table V lists the metadata we record for reproducibility. Datasheet peaks are not the sole inputs for validation: sustained bandwidth and tensor throughput from steady-state microbenchmarks (or conservative lower bounds) drive the generic roofline path; peak values are retained for upper-bound comparisons and for the stage-centric Blackwell validation kernels where appropriate. Hardware parameter files ship defaults for both primary platforms (B200, MI300A). We include a naive roofline baseline [10] (Troofline = max(FLOPs/Ppeak , bytes/BHBM )) as context, not as a competitive bar. Naive roofline uses only datasheet peaks and ignores cache hierarchies, pipeline stages, occupancy, and launch latency; it is not designed to predict execution time accurately. Table VI shows that roofline error exceeds 94% on all platforms, while our model achieves 1.3% (B200), 0.09% (MI300A), 4.7% (MI250X), and 9.6% (H200) on microbenchmarks. The gap illustrates why architecture-specific modeling is necessary on modern GPUs. TABLE VI: Microbenchmark validation: model MAE (%) per platform. Naive roofline shown as context (datasheet peaks only). Platform

Model MAE (%)

Roofline (%)

B200 MI300A H200 MI250X

21 27 21 19

1.33 0.09 9.57 4.69

96.1 99.6 94.5 97.9

B. Methodology and Validation Setup We validate both models against measured performance using optimized libraries (cuBLAS/CUTLASS, rocBLAS) and

hardware counters (Nsight Compute [28], ROCProfiler [29]). Each kernel runs 100 times after 10 warm-ups; we report median execution time (IQR of per-kernel error below 0.5% on MI300A, below 2% on B200). All reported MAE values use the base model (MWP=CWP=0, LDS terms not set). Table VII maps key model parameters to measured values. Table VIII gives the evaluation platform. Table IX summarizes validation workload classes. Hardware parameter files distinguish peak (datasheet) from sustained (microbenchmark) values for all platforms. a) Rodinia multi-segment modeling.: We model each Rodinia benchmark as a sum of segments (dominant GPU kernels or repeated launch patterns), each characterized by FLOPs, bytes, class, and an execution count nexec . Architecture-aware routing maps each segment class to the appropriate validated kernel family (stencil→transpose, compute-bound→GEMM, memory-bound→vector copy; Section IV). Measured time definitions. NVIDIA: sum of CUDA kernel durations from Nsight cuda_gpu_kern_sum. AMD: sum of HIP kernel durations from rocprof --stats. All reported MAE values use the same definition consistently within a platform. b) Segment construction and calibration.: Segment files are refined so that routing matches workload physics: • HotSpot (hs_calc): stencil class → memory-bound transpose proxy for grid traffic. • Pathfinder (dynproc_kernel): reduced effective FLOPs/bytes per step; effective timestep count aligned with profilers. • SRAD: single aggregate (N =M =0), traffic sized from bytes column. • Backprop: two layers merged into one compute segment to avoid double-counting launch latency. • Streamcluster: nexec scaled to measured launch regime. Table X reports MAE with default multipliers mcase =1. If residual error remains on new platforms, an optional calibration step can fit mcase with train/holdout separation. c) Blackwell B200.: Overall 1.31% MAE across 21 kernels on real B200 hardware (100 repeats, 10 warm-ups). • Memory-bound (vector add/copy/transpose/reduction): 8.4% MAE; vector ops 7–9% error from L2 benefits and 5–12 µs launch overhead. • Compute-bound (FP16/FP8/LLM GEMM via cuBLAS): 5.4% MAE; TMEM at 22 TB/s is conservative (24–26 TB/s in tuned kernels reduces error to 2–3%).

TABLE VII: Key model parameters and measurement method. Parameter

Value (or range)

Measurement method

TMEM BW (read/write) TMA latency LT M A tcgen05.mma latency SM Tensor throughput RT C Lmbar , Lcommit Infinity Cache / HBM BW L1/L2/LLC/HBM latency MI300A multi-kernel τinterf MI300A multi-GPU τinterf gpu MI300A tile selection MI300A fusion τfusion H2D/D2H effective BW Beff τmemcpy , τsync

16/8 TB/s 420 cyc 11–14 cyc 44.8–7702 TFLOPS 40–50 cyc 17.2 / 5.3 TB/s 5/50/150/400 cyc 50 µs (tuned) (accuracy) (tuned) 45 GB/s (default) 2 µs, 3 µs (defaults)

Microbenchmark: tile copy TMEM↔SMEM Microbenchmark: TMA copy latency Instruction timing (FP64–FP4) Peak sweep by precision Barrier microbenchmark Bandwidth microbenchmark Cache latency microbenchmark Tuned from concurrent-stream microbenchmark (1 vs. 2 streams) Tuned from multi-device microbenchmark (1 vs. 2 devices) Measured across GEMM tile sizes (8/16/32/64); predicted vs. measured best tile Tuned from unfused vs. fused GEMM+bias microbenchmark host_device_memcpy_bench_cuda / host_device_memcpy_bench_hip ANALYTICAL_T_MEMCPY_LAUNCH_S, ANALYTICAL_T_SYNC_S

TABLE VIII: Evaluation platforms used for validation. Item

Value

NVIDIA GPU AMD GPU Driver / CUDA / ROCm OS

NVIDIA Blackwell B200 (primary); NVIDIA H200 141 GB HBM3e (Rodinia / SPEChpc Tiny) AMD Instinct MI300A Accelerator (APU) ROCm 7.2.0 (MI300A); driver 550.163, CUDA 12.4 (B200/H200) Linux 4.18.0-553.16.1.el8 10.x86 64 (MI300A); Linux 6.x (B200/H200)

TABLE IX: Validation workload classes and benchmark suites. Class

Kernels

Suites

Memory-bound Compute-bound Balanced/Stencil

Vector add/copy/transpose, reduction FP64/FP16/FP8 GEMM (cuBLAS/rocBLAS) FFT, SpMV, GEMV; HotSpot stencil

Microbench, Rodinia (BFS, SRAD, Streamcluster), SPEChpc (505.lbm t, 518.tealeaf t) Microbench, Rodinia (Backprop), SPEChpc (521.miniswp t) Microbench, Rodinia (HotSpot, Pathfinder), SPEChpc (513.soma t, 532.sph exa t)

Balanced (FFT, SpMV, GEMV): 7.9% MAE; SpMV at 0.1% density shows 13.6% error (atomics, load balance not modeled). • 2-SM cooperative: predicted 1.30× speedup vs. measured 1.28× (within 2%). •

d) MI300A.: Overall ∼0.09% MAE across 27 kernels (vectors, reductions, 2D transposes, FP64 rocblas_dgemm, occupancy-tile GEMMs, VGPR/cache stencil variants). The roofline branch alone is optimistic for large 2D transpose traffic and sustained FP64 GEMM vs. peak MFMA; the model therefore applies host-measured calibration: 2 2 • Multipliers for 8192 and 16384 transposes. • Piecewise scaling vs. M =N =K for gemm_fp64. • Per-tile factors for 8×8 and 16×16 occupancy GEMMs (ordering preserved: 16×16 faster). Naive roofline error on MI300A remains ∼99% (Table VI). e) MI250X and H200.: MI250X: 4.7% MAE across 19 kernels (memory-bound vectors, FP64 GEMM, occupancy/tile study). FP64 GEMM tracks closely (e.g., 0.283 s predicted vs. 0.283 s measured at 163843 ). Tile ordering reproduced (16×16 faster). H200: We use the same model framework as B200 with updated HBM bandwidth (4.8 TB/s) and capacity (141 GB). Rodinia and SPEChpc validation results appear in Tables X and XI. Both H200 and MI250x use the same model frameworks as their primary counterparts with updated parameter files only.

TABLE X: Rodinia 3.1 validation: per-benchmark MAE (%) on B200 and MI300A. Benchmark

Class

B200

MI300A

hotspot_1024 hotspot_512 bfs_1M backprop_65536 pathfinder_1000 srad_502 streamcluster_1M

stencil stencil mem compute balanced balanced mem

31.0 15.4 44.9 33.0 0.4 0.5 12.4

23.6 1.6 40.9 21.3 0.1 0.5 0.03

C. Case Study 1: Rodinia 3.1 — Cross-Platform Both Rodinia [8] and SPEChpc [9] use the same pipeline: decompose each benchmark into kernel segments, route each to the appropriate model path (Section IV), predict total time, and compare against profiled GPU kernel sums. Tables X and XI report per-benchmark MAE on the primary platforms. Figure 4 compares predicted and measured execution time per Rodinia benchmark on both primary platforms. On B200, regular workloads (pathfinder 0.4%, srad 0.5%) are wellpredicted while irregular access (bfs 44.9%) and stencil (hotspot 31%) show higher error. On MI300A, the calibrated model achieves near-zero error on most benchmarks. For context, a naive roofline predictor on the same benchmarks yields ∼100% MAE overall on MI300A: for example, streamcluster_1M measures 157 ms but roofline predicts 0.005 ms (100% error), while our model predicts 157 ms (0.03% error). Even on the challenging bfs_1M, our 40.9% error compares to roofline’s 95.4%.

B200 Rodinia

B200 SPEChpc Tiny

MI300A Rodinia Measured Predicted

Measured Predicted

0.0%

MI300A SPEChpc Tiny Measured Predicted

10,000

Measured Predicted 1.5%

100

31.0%

0.4%

44.9%

0.5% 0.1% 1.3%

1.6%

0.0%

15.4%

0.10

1,000

18.5% 15.0%

0.2%

100

0.0%

32.7%

1.6%

0.1%

Execution time (ms)

0.5%

Execution time (ms)

12.4%

1.3%

0.0%

4 102 pot_

hots

512 pot_

6 553 p_6

bfs_1

pro

back

0.1%

33.0%

hots

00 er_1

find

path

_50

srad

amcl

stre

r_1M

uste

102

pot_

hots

512

pot_

hots

bfs_1

p_6

pro

back

553

path

0.8%

0.6%

0.8%

000

er_1

find

r_1M

_50

srad

uste amcl

_lbm

505

stre

Fig. 4: Rodinia: measured vs. predicted execution time per benchmark on B200 (left) and MI300A (right), log scale in ms. Per-benchmark MAE (%) annotated.

B200

MI300A

505.lbm_t 513.soma_t 518.tealeaf_t 519.clvleaf_t 521.miniswp_t 528.pot3d_t 532.sph_exa_t 534.hpgmgfv_t

mem bal mem mem comp mem bal mem

14.9 0.3 0.2 18.5 32.8 — 0.03 0.3

0.1 1.3 1.6 1.5 0.8 7.0 0.6 0.8

Table XI reports SPEChpc Tiny MAE; 535.weather_t is omitted for MI300A (no GPU kernels in profiler output). Figure 5 compares predicted and measured execution time per SPEChpc benchmark on both primary platforms. On B200, four benchmarks are under 1% MAE (soma 0.3%, tealeaf 0.2%, sph exa 0.03%, hpgmgfv 0.3%) with an overall MAE of 9.6%. On MI300A, all benchmarks are within 7% (profilerderived characterization; see Observation 3 in Discussion). E. Applying Analytical models to H200 and MI250x H200 and MI250X application results. When we apply the B200 model framework to H200 (parameter update only, no re-calibration) and the MI300A model to MI250X, applicationlevel MAE is higher: H200 Rodinia 43.6%, MI250X Rodinia 92.0%, H200 SPEChpc 555%, MI250X SPEChpc 59.3%. This is expected: the segment characterization (FLOPs, bytes) was derived on the B200 and MI300A and not on H200 and MI250X, so the model overpredicts or underpredicts due to bandwidth and cache hierarchy differences. Per-benchmark results for H200 and MI250X are available in supplementary materials. Summary. On the primary platforms, B200 achieves 1.3% and MI300A 0.09% microbenchmark MAE; MI300A Rodinia 12.5% and SPEChpc 1.3%. H200 and MI250x show higher application MAE consistent with cross-platform characterization mismatch. If MAE is below ∼5%, the model accurately captures the architecture; between 20–40%, the model provides useful estimates; above that, platform-specific recharacterization is needed.

iswp

_min

521

xa h_e

_sp

532

_hp

534

gmg

_lbm

505

_so

513

eale

518

le _clv

519

iswp

_min

521

xa h_e

_sp

532

fv gmg _hp

534

Model MAE by Platform and Workload Microbenchmarks Rodinia SPEChpc

102 10%

101

100

10 1

B200

D. Case Study 2: SPEChpc 2021 Tiny — Cross-Platform

le _clv

519

103

MAE (%)

Class

eale

518

Fig. 5: SPEChpc Tiny: measured vs. predicted execution time per benchmark on B200 (left) and MI300A (right), log scale in ms. Per-benchmark MAE (%) annotated.

TABLE XI: SPEChpc 2021 Tiny validation: per-benchmark MAE (%) on B200 and MI300A. 535.weather_t omitted (no GPU kernels in profiler). Benchmark

_so

513

H200

MI300A

MI250X

Fig. 6: Model MAE (%) by platform and benchmark suite. Microbenchmark MAE shown for all four platforms; application MAE shown where measured. VI. O BSERVATIONS Observation 1: Calibration and the MI300A 0.09% MAE. We observe systematic over- and under-prediction by workload class. On MI300A, the uncalibrated analytical model overestimates sustained bandwidth for large 2D transposes and peak MFMA throughput for FP64 GEMM. Hostmeasured calibration multipliers correct these cases, bringing the 27-kernel suite to ∼0.09% MAE. We note that this is a calibrated result: without per-case multipliers, the MI300A model achieves roughly 5–8% MAE, comparable to B200’s uncalibrated 1.31%. We report both because they serve different purposes: the uncalibrated model gives rapid, no-tuning estimates; the calibrated model shows ceiling accuracy when platform-specific effort is invested. We find that optional MWP/CWP limits and LDS bank-conflict terms can provide further improvement on specific kernels but are not required for the reported results. Observation 2: Per-benchmark Rodinia error analysis. The 12.5% MAE on MI300A is driven by three specific architectural outliers: bfs_1M (40.8% error): Irregular pointerchasing violates the model’s regular-access assumptions, as bandwidth depends on graph structure rather than working-set size. hotspot_1024 (23.7% error): A Stencil kernel where the transpose-proxy routing overestimates data reuse; the stencil pattern fails to map accurately to the 2D transpose-proxy bandwidth model. backprop_65536 (21.0% error): A small compute kernel where launch overhead and host–device syn-

TABLE XII: SPEChpc MI300A: profiler-derived vs. firstprinciples (FP) characterization. FLOP ratio = FP FLOPs / profiler FLOPs. Benchmark 505.lbm_t 513.soma_t 518.tealeaf_t 519.clvleaf_t 521.miniswp_t 528.pot3d_t 532.sph_exa_t 534.hpgmgfv_t

Prof. MAE

FP MAE

FLOP ratio

0.1% 1.3% 1.6% 1.5% 0.8% 7.0% 0.6% 0.8%

98.7% 31.8% 98.4% 98.7% 99.2% 10.3% 94.0% 19.4%

0.121 1.065 0.008 0.013 0.001 0.961 0.021 0.800

chronization dominate actual GPU execution time; the model captures compute and memory but underweights fixed-cost launch latency for microsecond-scale kernels. In contrast, pathfinder (0.1%), srad (0.5%), and streamcluster (0.03%) are regular memory-bound or balanced workloads where the wavefront model’s bandwidth and occupancy terms closely match measured behavior. We find this error distribution consistent across platforms: benchmarks with irregular access or very short kernels show higher error, while regular data-parallel workloads remain below 1%. This distribution suggests that the model’s accuracy boundary is defined by workload regularity rather than platform-specific artifacts. Observation 3: SPEChpc accuracy depends on characterization source. We observe that MI300A SPEChpc MAE is 1.3% when segment FLOPs/bytes come from profiler counters, but 92.5% when derived from first-principles algorithm analysis (source code, grid dimensions, stencil widths). To quantify this, we independently derived FLOPs and bytes for each SPEChpc benchmark from the source and ran both characterizations through the same model. The FLOP ratios (first-principles / profiler) reveal the gap: 521.miniswp_t has a ratio of 0.001 (5.1×109 vs. 4.8×1012 ), meaning the profiler-derived value is ∼1000× the source-code flop count. Similarly, 518.tealeaf_t (0.008) and 519.clvleaf_t (0.013) show orders-of-magnitude discrepancy. The two closest cases are 528.pot3d_t (ratio 0.96, first-principles MAE 10.3%) and 534.hpgmgfv_t (ratio 0.80, MAE 19.4%), where algorithm-level and profiler-level FLOP accounting roughly agree. Table XII summarizes the per-benchmark comparison. We emphasize that the 92.5% first-principles error is not a failure of the analytical model; it is a failure of the inputs. The model’s architecture-specific terms (Infinity Cache hierarchy, wavefront occupancy, VGPR pressure) are what reduce error from the naive roofline’s ∼206% to 1.3% on the same measured times. That 205× improvement is entirely due to the model correctly capturing MI300A’s execution behavior. The remaining gap between 1.3% (profiler-characterized) and 92.5% (first-principles) is a workload characterization problem, not a modeling problem: for OpenACC/OpenMP offload codes, the compiler generates GPU kernels whose actual FLOPs and memory traffic differ from source-level algorithm analysis by up to 1000× (e.g., 521.miniswp_t).

This compiler-generated kernel gap is orthogonal to the analytical model and represents an independent finding. We report profiler-derived MAE in the main tables (Table XI) because it isolates model accuracy from characterization accuracy, and we report the first-principles comparison here to quantify the characterization challenge for the community. Observation 4: H200 portability error analysis. The H200 results (43.6% Rodinia, 555% SPEChpc) are intentionally presented without re-calibration to show raw portability behavior. We observe two distinct failure modes. For Rodinia (43.6% overall), the B200-derived segment metadata transfers to H200 with moderate error. The worst case is hotspot_512 (88%): the segment file assumes a kernel decomposition tuned for Blackwell, but H200’s Hopper SM executes fewer, longer kernels with different memory traffic per launch. Regular benchmarks (srad 32.9%, pathfinder 42.1%) show that the Hopper roofline path captures the right order of magnitude but consistently overpredicts by 1.3–1.9×, suggesting the sustained bandwidth parameter needs H200-specific measurement rather than the B200 default. Notably, naive roofline is ∼100% on these same benchmarks, so even the uncalibrated port is substantially better. For SPEChpc (555% overall), the error is much larger because the segment FLOPs/bytes were characterized on MI300A (profiler-derived), not on H200. The MI300A characterization assumes Infinity Cache bandwidth (∼10.7 TB/s effective), but H200 sees only HBM bandwidth (∼4.2 TB/s sustained). This 2.5× mismatch in effective bandwidth propagates directly into predicted time: the model predicts MI300Ascale runtime on hardware that is 2–5× slower for these memory-bound codes. The exception is 521.miniswp_t (11% error), which is compute-bound; its prediction depends on tensor-core throughput, which scales more predictably across platforms than memory hierarchy behavior. We observe that naive roofline is uniformly ∼150% on H200 SPEChpc, meaning our uncalibrated model is actually worse than roofline for memory-bound SPEChpc on H200. This confirms that segment characterization must be platform-specific for memorybound workloads; compute-bound workloads transfer more reliably. Observation 5: Architectural differences in execution bottlenecks. We observe that Blackwell’s TMEM and decompression favor high-AI dense workloads (AI > 16 FLOPs/Byte), and that the explicit pipeline stages (TMA→TMEM→TC→Sync) lend themselves naturally to modular, stage-centric modeling. On MI300A, we find that balanced precision support (FP64–FP8) and unified physical memory fit heterogeneous CPU–GPU workloads, but that reaching the compute-bound regime requires higher reuse (AI > 23 FLOPs/Byte) than on Blackwell. The 256 MB Infinity Cache bridges this gap, delivering 1.5–2× over HBMbound bandwidth when workloads fit. We observe that roofline underestimates MI300A by 20–30% in these cache-resident cases precisely because it assumes a single HBM bandwidth. The ∼45% difference in AI thresholds between the two archi-

tectures suggests that tiling and autotuning strategies should be architecture-specific rather than addressing portability. Observation 6: Future of analytical models and adaptability to new GPUs. Every model coefficient maps to a microbenchmark, so adapting to a new GPU requires re-measuring parameters, not re-deriving formulas. The B200/H200 pair and MI300A/MI250X pair demonstrate this: same model framework, different parameter files. The model structure is determined by how the architecture accumulates results: dedicated TMEM (Blackwell stage model) vs. VGPR accumulators (CDNA wavefront model). For future GPUs within these families (Rubin, CDNA4), we expect parameteronly updates to suffice; a fundamentally new accumulation mechanism would require one new stage term. In our results, MAE below 15% after parameter update suggests the framework suffices; above 30% suggests a structural change. More broadly, our results suggest analytical models remain viable when architectures expose measurable execution phases, but the key bottleneck is shifting from model formulation to workload characterization: our first-principles experiment (Table XII) shows accurate FLOP/byte inputs are harder to obtain than accurate hardware parameters. As accelerators add more implicit hardware scheduling, hybrid analytical-ML approaches may be needed for the residual. Observation 7: Benchmark adequacy for modern accelerators. We observe that existing benchmark suites do not fully exercise the execution primitives that dominate modern GPU performance. Rodinia 3.1 [8], designed for early CUDA GPUs, does not use tensor cores, TMA, TMEM, or structured sparsity; its kernels are short (µs-scale) and exercise only basic memory and compute paths. SPEChpc 2021 Tiny provides more realistic HPC workloads but uses directive-based offload (OpenACC/OpenMP), which introduces a compiler-generated kernel layer between the algorithm and the hardware. We find that this compiler layer causes up to 1000× discrepancy between source-level and GPU-level FLOP counts (Table XII), making it difficult to attribute model error to the model vs. the characterization. Benchmark suites (a) representing native CUDA/HIP kernels that directly exercise tensor cores, TMA, and TMEM at representative problem sizes; (b) providing reference FLOP and byte counts derived from both algorithm analysis and hardware counters; and (c) spanning the full arithmetic intensity range from memory-bound streaming to compute-bound dense linear algebra, would be useful. Observation 8: Implications for hardware vendors. We observe several architecture-level differences with practical consequences for performance engineering. On the NVIDIA side, Blackwell’s dedicated TMEM (256 KB per SM) enables predictable, high-bandwidth accumulation that decouples tensor-core throughput from register pressure; this is a clear improvement over Hopper’s SMEM-based accumulators and contributes directly to the model’s low MAE. The TMA bulkcopy engine similarly reduces modeling complexity by making data movement explicit and measurable. We note, however, that the 2-SM cooperative execution model introduces a scheduling dependency that is harder to characterize: our 2-SM

predictions are within 2%, but only because the pairing is deterministic in current workloads. On the AMD side, MI300A’s 256 MB Infinity Cache is a significant architectural advantage for workloads with moderate reuse (working sets between 205–256 MB), delivering 1.5–2× effective bandwidth over HBM-only operation. However, we observe that the VGPRbased accumulation creates a tile-size vs. occupancy tradeoff that is harder to model than Blackwell’s TMEM approach: larger tiles improve cache reuse but reduce occupancy, and the optimal point depends on the specific VGPR allocation, which varies by compiler. Vendors could (a) expose sustained bandwidth and compute throughput in machine-readable format alongside datasheet peaks (our microbenchmarks had to measure what datasheets do not report); (b) provide deterministic kernel launch latency bounds (launch overhead dominates our error on short kernels like backprop_65536); and (c) for AMD, reduce sensitivity of occupancy to VGPR allocation through hardware register shadowing or compiler-managed spilling. Limitations. We note that both models assume regular compute and predictable memory access. In our experiments, accuracy degrades for irregular access patterns (sparse, indirection, atomics) and very short kernels where launch overhead dominates. Not yet modeled: cache replacement policy, coherence in multi-GPU configurations, thermal throttling, and power/energy. Based on our experience, the model supports: (1) procurement comparisons between B200 and MI300A without access to both; (2) autotuning guidance for tile size, occupancy, and precision; (3) co-design by identifying memory vs. compute bottlenecks; (4) rapid model instantiation on new hardware by running microbenchmarks. Future work. Future directions include power/energy models, hybrid analytical-ML for residual error, and applying the methodology to NVIDIA Rubin, AMD CDNA4, and Intel Gaudi3. VII. C ONCLUSION This paper presents analytical models for the current generation of GPU accelerators from NVIDIA and AMD: Blackwell B200 and CDNA3 MI300A. For Blackwell, we constructed a stage-centric model capturing TMEM, TMA, 5th-generation tensor cores, and the 2-SM cooperative execution model, to our knowledge the first validated execution-time model for this architecture. For CDNA3, we developed a wavefrontcentric formulation accounting for Infinity Cache hierarchy, VGPR register pressure, and occupancy-driven tile selection. Microbenchmark validation yields 1.31% MAE on B200 (21 kernels) and ∼0.09% on MI300A (27 kernels); naive roofline baselines exceed 95% error on the same kernels. Crossplatform validation against Rodinia 3.1 and SPEChpc 2021 Tiny confirms accuracy holds beyond the microbenchmark suite. The key architectural finding is that TMEM dominates execution time on matrix-heavy Blackwell kernels, which roofline cannot capture, while MI300A performance is gated

by occupancy and Infinity Cache reuse at arithmetic intensity thresholds roughly 45% higher than Blackwell. These structural differences require architecture-specific model terms, not just parameter substitution in a generic roofline. VIII. ACKNOWLEDGMENTS This material is based upon work supported by the U.S. Department of Energy, Office of Science, Office of Advanced Scientific Computing Research under prime contract DEACO5-000R22725 through the S4PST project funded by the Next Generation Scientific Software Technologies program. R EFERENCES [1] W. J. Dally, S. W. Keckler, and D. B. Kirk, “Evolution of the graphics processing unit (gpu),” IEEE Micro, vol. 41, no. 6, pp. 42–51, 2021. [Online]. Available: https://doi.org/10.1109/MM.2021.3113475 [2] K. Kurowski, M. Kulczewski, and M. Dobski, “Parallel and gpu based strategies for selected cfd and climate modeling models,” in Information Technologies in Environmental Engineering: New Trends and Challenges. Springer, 2011, pp. 735–747. [Online]. Available: https://doi.org/10.1007/978-3-642-19536-5 [3] N. Koilia and C. Kachris, “Hardware acceleration of llms: A comprehensive survey and comparison,” arXiv preprint arXiv:2409.03384, 2024, url = https://doi.org/10.48550/arXiv.2409.03384. [4] NVIDIA Corporation, NVIDIA Blackwell Architecture Technical Brief, NVIDIA, 2024. [Online]. Available: https://resources.nvidia.com/en-u s-blackwell-architecture [5] “Introducing amd cdna™ 3 architecture,” Advanced Micro Devices, Inc., Tech. Rep., 2023. [Online]. Available: https://www.amd.com/co ntent/dam/amd/en/documents/instinct-tech-docs/white-papers/amd-cdn a-3-white-paper.pdf [6] 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. [Online]. Available: https://doi.org/10.48550/a rXiv.2512.02189 [7] A. Jarmusch, C. Vitz, and S. Chandrasekaran, “Execution-centric characterization of FP8 matrix cores, asynchronous execution, and structured sparsity on AMD MI300A,” in Proceedings of the ACM International Symposium on High-Performance Parallel and Distributed Computing (HPDC). ACM, 2026. [Online]. Available: https://doi.org/10.48550/arXiv.2602.10262 [8] T. Lühnen, S. Nabi, and A. Koch, “Benchmarking thread block cluster,” in Proceedings of PDPTA. CSREA Press, url = https://doi.org/10.1109/HPEC62836.2024.10938416, 2024. [9] E. Gilman, H. Khaleghzadeh, K. Doshi, A. Dengel, B. Juurlink, N. Jain, and R. Ahmed, “Demystifying the placement policies of the NVIDIA GPU thread block scheduler for concurrent kernels,” ACM SIGMETRICS Performance Evaluation Review, vol. 48, no. 3, pp. 65–67, 2020. [Online]. Available: https://doi.org/10.1145/3453953.3453972 [10] S. Williams, A. Waterman, and D. Patterson, “Roofline: an insightful visual performance model for multicore architectures,” in Communications of the ACM, vol. 52, no. 4, 2009, url = https://doi.org/10.1145/1498765.1498785, pp. 65–76. [11] A. Ilic, F. Pratas, and L. Sousa, “Cache-aware roofline model: Upgrading the loft,” IEEE Computer Architecture Letters, vol. 13, no. 1, pp. 21–24, 2014. [Online]. Available: https://doi.org/10.1109/L-CA.2013.6 [12] G. Ofenbeck, R. Steinmann, V. C. Cabezas, D. G. Spampinato, and M. Püschel, “Applying the roofline model,” in 2014 IEEE International Symposium on Performance Analysis of Systems and Software (ISPASS). IEEE, 2014, pp. 76–85. [Online]. Available: https://doi.org/10.1109/ISPASS.2014.6842069 [13] M. Leinhauser, R. Widera, S. Bastrakov, A. Debus, M. Bussmann, and S. Chandrasekaran, “Metrics and design of an instruction roofline model for amd gpus,” ACM Transactions on Parallel Computing, 2022. [Online]. Available: https://doi.org/10.1145/3505285 [14] S. W. Hong and H. Kim, “An analytical model for a gpu architecture with memory-level and thread-level parallelism awareness,” in Proceedings of the 36th Annual International Symposium on Computer Architecture, 2009, url = https://doi.org/10.1145/1555754.1555775, pp. 152–163.

[15] J. Lee, H. Noh, S. Kim, J. Jeong, J. Kim, and J. Choi, “GCoM: A detailed GPU core model for accurate analytical modeling of modern GPUs,” in Proceedings of ISCA. ACM, 2022, pp. 424–436. [Online]. Available: https://doi.org/10.1145/3470496.3527384 [16] J.-C. Huang, J. H. Lee, H. Kim, and H.-H. S. Lee, “Gpumech: Gpu performance modeling technique based on interval analysis,” in MICRO. IEEE, 2014, pp. 68–79. [Online]. Available: https: //doi.org/10.1109/MICRO.2014.59 [17] L. Wang, M. Jahre, A. Adileh, and L. Eeckhout, “Mdm: The gpu memory divergence model,” in MICRO. IEEE, 2020, pp. 1009–1021. [Online]. Available: https://doi.org/10.1109/MICRO50266.2020.00089 [18] M. Khairy, Z. Shen, T. M. Aamodt, and T. G. Rogers, “Accel-Sim: An extensible simulation framework for validated GPU modeling,” in ISCA. IEEE, 2020, pp. 473–486. [Online]. Available: https: //doi.org/10.1109/ISCA45697.2020.00047 [19] 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). IEEE, url = https://doi.org/10.1109/ISPASS.2010.5452013, 2010, pp. 235–246. [20] Z. Jia, M. Maggioni, B. Staiger, and D. P. Scarpazza, “Dissecting the NVIDIA volta GPU architecture via microbenchmarking,” 2018. [Online]. Available: https://arxiv.org/abs/1804.06826 [21] W. Luo, R. Fan, Z. Li, D. Du, H. Liu, Q. Wang, and X. Chu, “Dissecting the NVIDIA hopper architecture through microbenchmarking and multiple level analysis,” 2025. [Online]. Available: https://arxiv.org/ab s/2501.12084 [22] J. Wahlgren, G. Schieffer, R. Shi, E. A. Leó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, 2025, arXiv:2508.12743, url = https://doi.org/10.1109/IISWC66894.2025.00038. [23] G. Schieffer, D. Medeiros, J. Faj, A. Marathe, and I. Peng, “Characterizing the performance, power efficiency, and programmability of AMD matrix cores,” in 2024 IEEE International Symposium on Performance Analysis of Systems and Software (ISPASS). Indianapolis, IN, url = https://doi.org/10.1109/ISPASS61541.2024.00022: IEEE, 2024. [24] R. Chowdhury, F. Silvestri, and F. Vella, “A computational model for tensor core units,” arXiv preprint arXiv:1908.06649, 2020, url = https://doi.org/10.48550/arXiv.1908.06649. [25] S. Cao, J. Wu, J. Chen, H. An, and Z. Yu, “Amali: An analytical model for accurately modeling llm inference on modern gpus,” in Proceedings of ISCA. ACM, 2025. [26] M. Fasi, N. J. Higham, M. Mikaitis, and S. Pranesh, “Numerical behavior of NVIDIA tensor cores,” PeerJ Computer Science, vol. 7, p. e330, 2021. [Online]. Available: https://doi.org/10.7717/peerj-cs.330 [27] X. Mei and X. Chu, “P-chase: A portable tool for measuring memory access characteristics on multicore computers,” in Embedded Software and Systems. Springer, 2009, pp. 76–83. [28] NVIDIA Corporation, NVIDIA Nsight Compute User Guide, 2025. [Online]. Available: https://docs.nvidia.com/nsight-compute/ [29] AMD ROCm Development Team, ROCProfiler: AMD ROCm GPU Profiling Tool, 2025. [Online]. Available: https://rocm.docs.amd.com/pr ojects/rocprofiler/en/latest/

Record · ID 168296 · SHA-256 71d747026a8ee54e
Retrieved via Conceptio — every document is proof-bundled with source, license, and retrieval metadata.