On the Limits of Performance Portability in Directive-Based GPU Programming Alessandro Romeo
arXiv:2606.12753v1 [cs.DC] 10 Jun 2026
SuperComputing Applications and Innovation Department CINECA Casalecchio di Reno, BO, Italy [email protected]
Nitin Shukla
SuperComputing Applications and Innovation Department CINECA Casalecchio di Reno, BO, Italy
Alessio Suriano
Andrea Mignone
Dipartimento di Fisica Università degli Studi di Torino Torino, Italy
Abstract The transition of scientific applications to GPU-accelerated exascale systems is constrained by trade-offs between performance, portability, and productivity. This work evaluates the performance portability of directive-based GPU programming by porting gPLUTO, a production-grade magnetohydrodynamics code for astrophysical simulations, from OpenACC to OpenMP, and analyzing its performance on NVIDIA A100 (Leonardo Booster) and AMD MI250X (LUMI-G) devices. On NVIDIA platforms, OpenACC and OpenMP achieve comparable performance due to a shared compiler backend, providing a consistent baseline for assessing algorithmic efficiency. In contrast, the same OpenMP implementation is approximately three times slower at the application level on AMD MI250X with respect to the NVIDIA A100 OpenACC baseline, with kernel-level slowdowns reaching up to an order of magnitude, driven by sensitivity to strided memory-access patterns and compiler limitations. Kernel-level profiling shows that the dominant contributors to runtime are memory-latency-bound rather than limited by peak bandwidth. In low-parallelism kernels, C++ abstraction layers increase register pressure and spilling, leading to extreme slowdowns of up to 47× in specific cases. These results indicate that portable performance across GPU architectures requires not only application-level changes but also continued advances in compiler backends and architecture-aware optimization strategies.
Keywords GPU offloading, OpenMP, OpenACC, performance portability, High performance computing, gPLUTO, GPU benchmarking, performance Pre-print version of the manuscript submitted to The 23rd ACM International Conference on Computing Frontiers (CF’26). Permission to make digital or hard copies of all or part of this work for personal or classroom use is granted without fee provided that copies are not made or distributed for profit or commercial advantage and that copies bear this notice and the full citation on the first page. Copyrights for components of this work owned by others than the author(s) must be honored. Abstracting with credit is permitted. To copy otherwise, or republish, to post on servers or to redistribute to lists, requires prior specific permission and/or a fee. Request permissions from [email protected]. Conference’17, Washington, DC, USA © 2026 Copyright held by the owner/author(s). Publication rights licensed to ACM. ACM ISBN 978-x-xxxx-xxxx-x/YYYY/MM https://doi.org/10.1145/nnnnnnn.nnnnnnn
Stefano Truzzi
Dipartimento di Fisica Università degli Studi di Torino Torino, Italy
Dipartimento di Fisica Università degli Studi di Torino Torino, Italy optimization, heterogeneous systems, NVIDIA A100, AMD MI250X, Leonardo, LUMI ACM Reference Format: Alessandro Romeo, Nitin Shukla, Stefano Truzzi, Alessio Suriano, and Andrea Mignone. 2026. On the Limits of Performance Portability in DirectiveBased GPU Programming. In . ACM, New York, NY, USA, 8 pages. https: //doi.org/10.1145/nnnnnnn.nnnnnnn
1
Introduction
High-performance computing systems have become increasingly heterogeneous, with GPUs driving progress toward exascale computing (1018 calculations per second): over 50% of TOP500 systems now integrate accelerators to maximize throughput and energy efficiency 1 . While GPUs have opened new frontiers in scientific simulations, legacy applications face significant portability challenges, reflecting the long-standing trade-off between performance, portability, and productivity [13, 20, 33]. Several EuroHPC Centers of Excellence across scientific domains - including SPACE2 [38, 39], Plasma-PEPSC3 [47], MaX4 [18], and ChEESE-2P5 [16] - support reengineering flagship codes for exascale performance and portable HPC best practices. Transitioning large legacy scientific codes to GPUs raises fundamental questions: (i) Can they be ported with minimal effort? (ii) How can reproducibility be preserved across evolving hardware? (iii) How can performance portability be achieved without sacrificing efficiency? Exploiting modern GPU memory bandwidths (2-3 TB/s) requires careful parallelism exposure and memory-access management. Lowlevel APIs (CUDA, HIP) promise peak performance but lack portability. C++ abstraction libraries (SYCL, RAJA, Kokkos) provide portability through templates but require extensive code modifications [8]. Directive-based models (OpenACC [31], OpenMP [32]) improve portability by abstracting device details, though often with performance penalties [22, 24, 27]: OpenACC benefits from a mature NVIDIA CUDA backend, while OpenMP offloading shows uneven optimization across compilers and architectures [4, 9]. While quantitative 1 https://www.top500.org/lists/top500/2025/11/ 2 https://www.space-coe.eu/ 3 https://plasma-pepsc.eu/
4 https://www.max-centre.eu/ 5 https://www.cheese2.eu/
Conference’17, July 2017, Washington, DC, USA
portability metrics exist [21, 25, 34, 37], applying them to multikernel production scientific applications remains challenging. This work focuses on directive-based models due to their widespread adoption in legacy MHD-Godunov [19, 42] and particlein-cell [28, 40] codes. Large-scale astrophysical and plasma codes exemplify how memory latency, compiler behavior, and data layout interact to determine achievable performance on modern accelerators. In particular, this paper examines performance portability limitations of directive-based GPU codes by analyzing the porting of gPLUTO 6 , a production-grade MHD code for astrophysical plasma simulations, from OpenACC to OpenMP. We hypothesize that such limitations arise primarily from compiler implementation differences rather than programming model semantics, though architectural characteristics (memory hierarchy organization and sensitivity to strided memory accesses) play a crucial role. To test this, we perform kernel-level analysis of throughput, arithmetic intensity, instruction-level parallelism, register usage, and cache efficiency across GPU architectures and compiler toolchains, exploring performance portability challenges and trade-offs for a real application memory-latency dominated [11, 12, 26, 44, 45]. The remainder of this paper is organized as follows. Section 2 describes the methodology and performance metrics used throughout the study, as well as the practical challenges encountered during the porting process. Section 3 presents application-level and kernellevel results for gPLUTO on NVIDIA A100 and AMD MI250X GPUs. Section 4 analyzes cross-architecture portability. Finally, Section 5 summarizes the main findings and discusses their implications for future GPU software development. This study operates under unified memory on both platforms (NVIDIA Managed Memory on A100, AMD Unified Memory on MI250X), as required for stable execution of gPLUTO’s pointer-intensive C++ codebase. While this introduces vendor-specific runtime behaviors (page migration, fault handling), our analysis focuses on programming model semantics and compiler-generated code; disentangling memory management from compiler and architectural effects remains future work.
2
Methodology and Experimental Setup
This section describes the experimental methodology adopted in this work. We present the targeted HPC clusters, as well as the production application gPLUTO, the adopted strategies implemented to optimize its OpenMP translation from OpenACC, the profiling tools, and the performance and execution metrics employed throughout the analysis.
2.1
Platform Details and Compilation Settings
Experiments were conducted on NVIDIA A100 (Leonardo Booster) and AMD MI250X (LUMI-G) using single CPU-GPU nodes to isolate accelerator performance and eliminate inter-node communication effects. Table 1 summarizes the architectural differences. While the A100 features a monolithic design [7], the MI250X is a multi-chip module (MCM) where each Graphics Compute Die (GCD) exposes distinct memory domains [41]. The MI250X offers higher peak FP64 throughput and larger register files, whereas the A100 provides significantly larger L2 capacity and a configurable L1/shared memory [5]. In particular, MI250X employs a disaggregated per-CU 6 https://plutocode.ph.unito.it/
Alessandro Romeo, Nitin Shukla, Stefano Truzzi, Alessio Suriano, and Andrea Mignone Table 1: GPU architecture specifications: Leonardo Booster (NVIDIA A100) and LUMI-G (AMD MI250X, single GCD). Work-items ≡ NVIDIA threads. Specification Compute Units Peak FP64 (TFLOP/s) Max Base Clock (GHz) HBM Bandwidth (TB/s) HBM Capacity (GB) L1 Cache (KB) L2 Cache (MB) Max 32-bit registers Vector register file [KB] Execution group size Register file scope
A100 108 SMs 9.7 1.41 2.0 64 192 40 255 (per thread) 256 (64K × 32-bit per SM) Warp = 32 threads per SM
MI250X (1 GCD) 110 CUs 13.3 (26.5 full MCM) 1.70 1.6 (3.2 full MCM) 64 (128 full MCM) 64 8 256 (per work-item) 512 (per CU) Wavefront = 64 threads per CU
front-end with smaller caches (16 KB vector L1 per CU, plus 16 KB scalar L1 and 32 KB instruction L1 per 2 CUs) [2]. Beyond an almost identical per-work-item cap (256 vs. 255), MI250X provisions a larger CU-level vector register file (≈512 KB per CU) than A100 per SM (64K 32-bit regs ≡ 256 KB per SM). Note that the register-file difference is independent of cache sizes [2, 29]. These architectural variations, summarized in Table 1, directly impact memory-latency tolerance and occupancy7 . On Leonardo (NVIDIA A100), both the OpenACC and the OpenMP versions are compiled with the NVIDIA HPC SDK (nvc++) 24.5, using managed memory and relocatable device code. Optimization and inlining are enabled to match production settings. On LUMI we use ROCm 6.2.4 and the LLVM/Clang OpenMP offload toolchain8 provided by the system environment with -O3 optimization. To ensure a fair comparison with the monolithic A100 and avoid variable NUMA placement effects, singleGPU runs on LUMI were explicitly bound to a specific GCD using ROCR_VISIBLE_DEVICES. This setting does not alter kernel execution behavior or enable additional optimizations. Both platforms utilized unified memory mechanisms (NVIDIA Managed Memory and AMD’s CRAY_ACC_USE_UNIFIED_MEM=1 and HSA_XNACK=1) to support gPLUTO’s pointer-intensive C++ structures. While transparent page migration may introduce performance overhead compared to explicit target data regions, these settings were necessary for code stability across both architectures.
2.2
gPLUTO: Production Astrophysics Code
gPLUTO is a finite-volume, Godunov-type code solving gas and plasma dynamics equations in astrophysical contexts. It supports five physics modules ranging from classical hydrodynamics to resistive relativistic magnetohydrodynamics. 2.2.1 Multi-Vendor GPU Implementation of gPLUTO. gPLUTO porting began in 2023 within SPACE-CoE [35], offloading approximately 70% of modules to GPUs with strong NVIDIA performance. The code was redesigned for exascale GPU computing, using OpenACC for acceleration and scaling efficiently to thousands of GPUs with non-blocking MPI and asynchronous data exchange. Development involved extensive parallelization optimization and a C-to-C++ rewrite introducing multidimensional array classes and templates. However, limited OpenACC compiler support for AMD C/C++ 7 Note that registers are allocated per thread/work-item (32-bit) while occupancy is
limited by the SM/CU-level register file and by the execution group granularity. 8 (i) craype-x86-trento, (ii) PrgEnv-amd/8.5.0, (iii) craype-accel-amd-gfx90a, and (iv) rocm/6.2.4 (clang++).
On the Limits of Performance Portability in Directive-Based GPU Programming
prompted an OpenMP adoption [43]. OpenMP directives were inserted manually or via Intel’s automated translation framework9 , though manual tuning remains necessary for performance-critical kernels. This approach preserves productivity while achieving vendor-neutral portability. 2.2.2 Practical Lessons: OpenACC to OpenMP Porting. In agreement with other works, achieving OpenACC-comparable OpenMP performance requires careful redesign beyond mechanical pragma replacement. Such a process involves refactoring, aliasing control, and careful handling of data abstractions [3, 17]. While OpenMP offloading has matured through multiple compiler generations [6, 10, 49], production deployment still requires detailed profiling and kernel-specific tuning. For the work within gPLUTO, key findings include: Parallelism exposure and nested loop restructuring: although OpenMP 5.x has relaxed perfect-nesting constraints, GPU offloading often exposes only two reliable parallelism levels (teams and threads), as simd support and its mapping to GPU lanes remains compiler-dependent. Optimal directives (teams loop vs. teams distribute parallel for) vary by kernel, compiler (NVHPC vs. ROCm), and architecture. Achieving full occupancy frequently requires collapse(3) on spatial loop nests, necessitating code restructuring: intermediate declarations must move inside innermost loops, improving exposed parallelism but increasing register pressure through repeated initialization and reduced compiler hoisting. In contrast, OpenACC’s gang/worker/vector hierarchy maps iteration spaces without full collapsing, allowing intermediate declarations outside vector loops and reducing register pressure. #pragma omp target teams distribute parallel for collapse(2) #pragma acc parallel loop collapse(2) for (k = kbeg; k <= kend; k++){ for (j = jbeg; j <= jend; j++){ // Intermediate declarations / temporaries not depending on i // (can stay here in OpenACC thanks to vector level) ... // No guaranteed omp simd mapping across compilers #pragma acc loop vector for (i = ibeg; i <= iend; i++){ ... } } }
Optimizations improving OpenMP performance may therefore degrade it through register pressure, while OpenACC achieves equivalent parallelism without restructuring: optimizing for one model may produce a performance loss for the other. Moreover, nonperfect loop collapsing availability remains limited and compilerdependent on EuroHPC clusters. Data management and C++ abstraction handling: while OpenMP 5.x supports implicit mapping and unified shared memory, optimal performance for pointer-based C++ structures requires explicit map clauses or declare mapper directives. Without them, compilers generate redundant private copies, increasing register pressure and spilling. OpenACC handles such abstractions more robustly via implicit deep-copy semantics. For instance, OpenACC’s present clause enforces device residency with runtime errors; OpenMP’s present modifier silently creates new mappings if data is absent, risking unexpected transfers. Complex structures require declare mapper or pointer APIs (e.g., use_device_ptr, 9 https://github.com/intel/intel-application-migration-tool-for-openacc-to-openmp
Conference’17, July 2017, Washington, DC, USA
omp_target_alloc), increasing complexity. Incorrect mapper definitions can cause silent data corruption or device segmentation faults, with limited compiler diagnostics to aid debugging. Additionally, mapper support varies across compilers, and complex nested structures or templated wrappers may prevent successful instantiation, increasing maintenance burden and significant code modifications. Compiler-specific issues: they include NVHPC 24.5 internal errors with separated teams distribute + parallel for (resolved in 24.7), incomplete simd support in ROCm 6.2.4, and subtle differences in reduction semantics, sequential loop enforcement and device function declarations. 2.2.3 gPLUTO’s main cycle and adopted tests. Concerning gPLUTO behavior, the code evolves conservation laws through three key stages — reconstruction of primitive (or characteristic) variables, Riemann problem solution at cell interfaces, and flux-conservative temporal update — executed as distinct GPU kernels repeated per spatial dimension. Additional procedures are required in the presence of magnetic fields (e.g., constrained transport routines) and for Runge-Kutta (3rd -order by default) time integration, as well as for setting boundary conditions and performing domain exchanges through MPI communication. We focus on the 3D Orszag–Tang vortex problem (3523 grid) and the 2D Riemann test, profiling kernels Reconstruct(), HLLD_Solver() (or HLL_Solver() for Riemann 2D), RightHandSide(), and auxiliary routines (CT_EMF(), US()). The 3D Orszag–Tang problem with 3523 grid (≈ 43.6 million cells) was selected to represent realistic gPLUTO production workloads while ensuring adequate GPU utilization. This configuration generates approximately 170,000 thread blocks (⌈3523 /256⌉), providing 1,500–1,600 blocks per SM/CU and exceeding hardware concurrency limits. Profiling confirms that observed occupancy limitations (12–71%, Tables 2 and 5) originate from register pressure and memory-latency effects rather than insufficient exposed parallelism. The 4.3× higher parallelism relative to the 2D Riemann case (32002 ≈ 10.2 million cells) enables both OpenACC and OpenMP to effectively hide register-spilling latency in 3D.
2.3
Profiling Tools and Performance Metrics
Performance analysis employed NVIDIA Nsight Compute/Systems for A100 and AMD rocProfiler/OmniPerf for MI250X. These tools collect hardware counters for throughput, bandwidth, cache efficiency, and instruction-level metrics following vendor roofline methodologies [1, 15, 30]. Metrics are derived from architecture-specific hardware counters. These metrics are functionally analogous across platforms but not numerically identical; cross-architecture comparisons therefore identify execution regimes qualitatively rather than comparing absolute values. We report achieved metrics (occupancy, bandwidth) measured at runtime rather than theoretical peaks, as they reflect real kernel execution under register pressure, memory latency, and instruction scheduling. Kernel runtimes (ms) and fraction per step (%): kernel execution times on NVIDIA A100 are obtained from Nsight Systems, while on AMD MI250X they are obtained from ROCm profiling tools (RocProf/OmniPerf). Reported single-kernel times correspond to steady-state execution and may slightly overestimate per-step
Conference’17, July 2017, Washington, DC, USA
Alessandro Romeo, Nitin Shukla, Stefano Truzzi, Alessio Suriano, and Andrea Mignone
costs due to initialization and profiling overheads. Only the dominant computational kernels are included in the runtime breakdown; auxiliary kernels (e.g., boundary conditions) account for the remaining execution time. The execution times reported in Section 3 refer to a single kernel invocation (e.g., one directional sweep for one Runge-Kutta stage) measured under isolated profiling conditions. Runtime fractions are computed per Runge-Kutta time step by aggregating all directional variants and stage repetitions of each kernel. Fractions are derived from application runs with profiling enabled; while absolute timings may be affected by profiling overhead, relative kernel contributions are preserved. The same methodology is applied consistently to both Leonardo and LUMI results. Computational Throughput (GFLOP/s): this metric quantifies sustained double-precision floating-point (FP64) operations per second: FP64 Computational Throughput =
FP64 FLOPs Runtime (s) × 109
.
It is derived from hardware counters normalized for vector width on AMD. VALU utilization on MI250X indicates instruction issue cycles, not peak throughput, making FP64 throughput the preferred metric. This choice is further justified a posteriori, as all examined gPLUTO kernels predominantly execute FP64 arithmetic operations. Achieved Memory Bandwidth (GB/s): data transferred per second between L2 cache and high-bandwidth memory (HBM), capturing DRAM traffic: Memory Bandwidth =
Bytes Transferred Runtime (s) × 109
.
This quantifies HBM pressure following the aforementioned hierarchical roofline methodology. High L2 hit rates reduce HBM bandwidth without reducing memory activity. Arithmetic Intensity (AI): FP64 FLOPs per HBM byte transferred: AI =
FP64 FLOPs
Bytes Transferred
.
HBM-based AI enables consistent identification of compute- versus memory-limited regimes across architectures [7, 36, 48] without reconstructing pre-cache traffic, which is architecture-dependent and not directly comparable. Additional Metrics: instructions per cycle (IPC, due to fundamental differences in execution models, used for intra-architecture analysis only); register usage (measuring spilling probability into local memory); achieved occupancy (the fraction of active warps or wavefronts relative to hardware limits, reflecting kernels’ ability to hide latency through concurrency); L1/L2 cache hit rates (to estimate spatial locality); branch utilization and efficiency (to quantify control-flow divergence within warps or wavefronts); memorylatency indicators. Kernel Classification: kernels are classified as compute-bound (high AI, high throughput), memory-bandwidth-bound (low AI, high achieved bandwidth), memory-latency-bound (low bandwidth despite frequent operations, low utilization), or control-flow-bound (branch-dominated with low AI, low execution-unit utilization, and limited memory throughput). Empirically a compute-bound kernel exhibits AI > 5 flop/byte and utilization > 30%; latency-bound exhibits AI < 2 flop/byte and bandwidth < 20% of peak.
3
Results
In the following sections OpenACC on NVHPC represents the most mature implementation and establishes performance upper bounds. Kernel metrics represent median or average values from steady-state time steps, excluding initialization overheads. Crossarchitecture comparisons use the same OpenMP implementation to isolate architectural effects.
3.1
gPLUTO Performance on Leonardo (NVIDIA A100)
Table 2 presents average metrics across directional kernel variants for the 3D Orszag-Tang case. Nsight Systems profiling reveals mild directional dependence: HLLD() 𝑗-sweeps execute approximately 20% slower than 𝑖, 𝑘 sweeps on A100, while RHS() exhibits approximately 14% directional variation. The US() kernels demonstrate factor-of-two runtime differences but maintain similar performance characteristics. Metrics are hence averaged across directional variants (𝑖/𝑗/𝑘) for the same kernel type and across multiple profiling runs to reduce measurement noise. HLLD() is primarily instruction-latency-bound: extreme register pressure (> 170/thread) collapses occupancy to 12% and triggers massive spilling (approximately 8.8 × 107 spill loads and 4.9 × 107 spill stores per launch), independently of the programming model used. This prevents the SM from hiding latency due to insufficient concurrent warps, despite high cache hit rates (L1 is approximately 50% while L2 is 87%). Low SM utilization (< 1%) and low throughput (approximately 50 GB/s versus 2 TB/s peak) decoupled from HBM traffic confirm that register-resident operations and latency, rather than bandwidth or memory locality, are the bottlenecks. In fact, HBM-based AI slightly underestimates sustained throughput. Even minor register increases disproportionately degrade performance, as SMs cannot hide instruction latency. Differences in branch instruction counts reflect compiler-level if-conversion and predication, but do not translate into runtime improvements due to the kernel being latency-bound. In contrast, Reconstruct() achieves balanced performance with moderate AI (≈ 6 flop/byte), acceptable occupancy (37–48%), and regular access patterns. L1 locality is high (approximately 63%), and L2 reuse is also high (approximately 78%). No register spilling occurs with 64–74 registers per thread. The near-identical performance between OpenACC and OpenMP confirms the efficiency of the unified CUDA backend for kernels with low register pressure. RHS(), CT_EMF(), and US() remain memory-latency-bound (AI < 0.4 flop/byte). Despite 37-64% occupancy, all exhibit limited L1 locality (18–57%) despite moderateto-high L2 hit rates (56–84%), minimal spilling, yet bandwidth 2-3 orders below peak. CT_EMF() bandwidth variability reflects directional access anisotropy10 , not improved computational efficiency. CT_EMF() shows the largest OpenACC-OpenMP gap, suggesting its OpenMP mapping/privatization is suboptimal; kernel-specific tuning (loop restructuring, temporary placement) would likely improve parity11 . 10 This term describes performance variation when the same kernel accesses multidi-
mensional arrays along different spatial dimensions.
11 CT_EMF() bandwidth varies directionally (OpenACC: 83 GB/s avg; OpenMP:
181–418 GB/s per sweep), yet execution time differs by <10%, confirming latencybound behavior.
On the Limits of Performance Portability in Directive-Based GPU Programming
Conference’17, July 2017, Washington, DC, USA
Table 2: gPLUTO kernel metrics on NVIDIA A100 (NVHPC 24.5), 3D Orszag-Tang. OpenACC and OpenMP achieve near-identical performance (except CT_EMF()) via the unified CUDA backend. Metric Time (ms) Runtime fraction (%) FP64 ThrP (GFlop/s) Mem BW (GB/s) AI (flop/byte) Occupancy (%) IPC L1 Hit (%) L2 Hit (%) Registers/thread Spill Loads ×107 Spill Stores ×107 Branch Instr. ×107 Avg Divergent Branches Branch Efficiency (%)
HLLD() ACC OMP 31.6 32.0 20.5 16.7 71.3 66.9 48.9 42.1 1.46 1.61 12 12 0.10 0.10 50.4 58.1 87.5 87.2 172 182 8.8 9.0 4.9 5.3 15.8 0.4 501 450 99.7 99.8
Reconstruct() ACC OMP 13.7 14.6 8.9 7.6 217.1 213.6 35.3 35.0 6.16 6.12 48 37 0.51 0.54 62.7 62.7 78.1 77.9 64 74 0 0 0 0 37.4 10.8 0 0 100 100
Overall, no kernels reach compute- or bandwidth-bound regimes. OpenMP exhibits a systematic 5-10% slowdown due to slightly higher register pressure, necessitating iterative tuning to reach near-parity with OpenACC. 3.1.1 OpenMP Fragility in Low-Parallelism Regimes: 2D Riemann. The 2D Riemann test (3200 × 3200 × 1 domain, over 107 iterations) reveals stark OpenACC-OpenMP differences absent in 3D configurations. Table 3 shows that OpenMP is 47× slower for Reconstruct() and 6× slower for the HLL() Riemann Solver, despite a nominally large iteration space. Unlike many small 2D test cases reported in Table 3: Riemann 2D kernel metrics on NVIDIA A100. OpenMP exhibits severe register spilling and suboptimal grid mapping despite large iteration space, causing order-of-magnitude slowdowns. Metric Grid size Time (ms) FP64 ThrP (GFlop/s) Mem BW (GB/s) Memory Throughput (%) AI (flop/byte) Occupancy (%) L1 Hit (%) L2 Hit (%) Registers/thread Spill Loads ×107 Spill Stores ×107 Branch Instr. ×107 Avg Divergent Branches Branch Efficiency (%)
Rec_ACC() 3200 × 128 1.24 17.9 1003 63.03 1.29 20.14 66 76 104 0 0 1.2 67.82 99.53
Rec_OMP() 1257 × 128 58 24.55 240 14.76 0.15 24.25 61 88 98 8.8 5.5 8.8 119.5 99.92
HLL_ACC() 3200 × 128 1.63 15.06 1000 61.24 1.07 30.28 10.87 66 90 0 0 0.69 19.35 99.77
HLL_OMP() 3200 × 128 9.84 27.73 505 31 0.38 33.8 66.5 79 80 4.4 1.1 1.8 19.35 99.93
the literature [10, 23, 46], this poor performance cannot be attributed to insufficient parallelism. Two root causes explain this behavior. First, register spilling: OpenMP generates 8.8×107 spill loads for Reconstruct() (compared with zero for OpenACC), dominating execution. This spilling is driven by conservative privatization of C++ wrapper objects (Ary1D/Ary2D) without using mappers rather than by the raw register count, which is actually comparable across models. Second, suboptimal grid mapping: NVHPC OpenMP selects 1257 × 128 grid vs. 3200 × 128 for OpenACC, exposing 2.5× less parallelism and reducing latency hiding capability. This compilerdriven grid selection does not reflect OpenMP model limitations; enforcing larger grids requires num_teams or thread_limit clauses, introducing machine-specific tuning that undermines portability. The trivial (𝑘 = 1) dimension exacerbates this; OpenACC handles it transparently through automatic collapse, while OpenMP exhibits sensitivity to local object placement and nested loop structure.
RHS() ACC OMP 9.0 9.3 5.8 4.9 11.0 11.0 32.1 31.0 0.34 0.36 54 48 0.08 0.09 57.5 49.7 79.8 79.5 56 64 0 0 0 0 1.8 0.1 0 0 100 100
CT_EMF() ACC OMP 7.8 23.9 10.1 25.0 15.5 18.7 83 278 0.19 0.10 37 46 0.02 0.02 18.3 18.5 68.2 55.8 80 65 0 0 0 0 0.9 0.14 502 0 95.5 100
US() ACC OMP 18.8 16.6 12.2 8.7 0.85 0.76 22.7 22.1 0.04 0.05 64 52 0.09 0.14 23.5 26.6 81.7 84.4 38 56 0 0 0 0 12 15.2 0 0 100 100
To quantify this sensitivity, we considered the best-case OpenMP configuration, analogous to current OpenACC implementation, which maps iteration space to GPUs. This is shown in the following snippet: #pragma acc parallel loop collapse(2) present(d, Dts, grid) #pragma omp target teams loop collapse(2) for (k = kbeg; k <= kend; k++){ for (j = jbeg; j <= jend; j++){ long int offset = ni*(j + nj*k); Ary1D cmax(&Dts->cmax[offset],ni); ... Ary2D vL(&d->sweep.vL[offset1],ni,NVAR); ... #pragma acc loop #pragma omp loop for (i = ibeg; i <= iend; i++){ ... } } }
and was compared against three alternative offloading strategies: (i) omp target teams loop collapse(3), with all declarations inside the i-loop; (ii) omp target teams distribute parallel for collapse(3), with all declarations moved inside the innermost i-loop; (iii) omp target teams distribute parallel for collapse(2), leaving declarations between the i- and j-loops and applying omp simd before the i-loop. Table 4: OpenMP directive sensitivity for HLL() (Riemann 2D): minor syntax changes cause up to 22 × runtime variation. Metric Grid size Time (ms) Occupancy (%) Registers/thread Local Mem (GB) Dynamic Shared Memory (bytes) Latency (𝜇 s)
Best 3200 × 128 9.84 37.5 80 0.34 192 5.865
Case 1 80025 × 1 11.7 43.75 72 0.44 0 6.239
Case 2 80025 × 1 20.2 25 114 0.51 0 6.246
Case 3 25 × 1 217 25 112 0.51 0 6.985
Table 4 quantifies OpenMP directive sensitivity for the HLL() kernel. Alternative formulations vary runtime by 22× through grid serialization and increased spilling. Only the best-performing variant allocates shared memory (192 bytes), reducing local-memory traffic. Directive changes inhibit this optimization. A final example of OpenMP rigidity emerges when applying collapse to the Reconstruct() kernel. This kernel contains several GPU-managed array wrappers (Ary1D, Ary2D) and local matrices with macrodefined dimensions that are not always statically known to the
Conference’17, July 2017, Washington, DC, USA
Alessandro Romeo, Nitin Shukla, Stefano Truzzi, Alessio Suriano, and Andrea Mignone
compiler. These features prevent NVHPC 24.5 from proving iteration independence across all three loops. Both collapse(3) and teams distribute parallel for fail as a result. Simplifying these abstractions into plain C-style pointer arrays can restore parallelization and reduce spilling. However, this requires intrusive refactoring and breaks one-to-one correspondence with the OpenACC version. In summary, high parallelism in the 3D Orszag–Tang problem enables both models to hide spilling latency. In 2D, limited concurrency combined with C++ abstraction overhead causes OpenMP performance breakdown. This confirms performance portability between OpenACC and OpenMP is strongly kernel-size-dependent, with small-to-medium kernels constituting the most challenging regime for OpenMP GPU offloading.
latency under irregular access patterns. Extremely low AI and VALU utilization are observed. Execution is dominated by stalled memory operations rather than sustained data streaming. Unlike the A100, MI250X kernels exhibit moderate VGPR usage (56-120) and negligible spilling, though scratch allocation is non-zero. Slower HLLD 𝑗,𝑘 () uses fewer VGPRs (108) than HLLD𝑖 () (120). Register pressure is not the limiting factor. Although the per-work-item VGPR cap is nearly identical (MI250X 256 vs. A100 255), the larger CU-level register file on CDNA2 mitigates abstraction-induced spilling, shifting the primary bottleneck to memory-access patterns and unified-memory runtime page migration effects exacerbated by irregular or strided accesses [41].
3.2
Cross-architecture analysis reveals substantial MI250X performance deterioration despite identical OpenMP code, as shown by total execution times in Figure 1. While both architectures sustain non-negligible HBM traffic, MI250X converts it into lower FP64 throughput due to poor memory coalescing under strided access. The much smaller L1/L2 cache (16 KB vector L1 per CU, 8 MB L2 vs. 192 KB L1 and 40 MB L2 on A100) amplifies non-unit-stride penalties, consistently with the low hit rates and VMEM underutilization in Table 5. The HLLD() 3.7× directional anisotropy on MI250X (vs. ≈20% on A100) exemplifies this: strided 𝑗, 𝑘 sweeps bypass L1 with hit rate drops, collapsing VALU utilization and increasing VMEM stalls. Application-level performance is dictated by these latency-sensitive kernels: large FLOP/s gains (16.5× vs. A100) and high VALU utilization in Reconstruct() do not compensate previous latencies, resulting in a 3× slowdown shown in Figure 1, illustrating why aggregate throughput metrics obscure real bottlenecks.12
Table 5 reports MI250X obtained results. One can suddenly note that performance differs from A100, despite identical code. Table 5: gPLUTO kernel metrics on AMD MI250X (ROCm 6.2.4), 3D OrszagTang. Metrics reflect single-GCD execution; aggregate MI250X peak (220 CUs, 3.2 TB/s) represents the full MCM capability. HLLD𝑖 () benefits from unit-stride access; other kernels remain latency-bound. Metric Time (ms) Runtime fraction (%) FP64 ThrP (GFlop/s) Mem BW (GB/s) AI (flop/byte) Occupancy (%) IPC L1 Hit (%) L2 Hit (%) VGPR/wavefront Scratch Alloc. (bytes) VALU Utilization (%) VMEM Utilization (%) Branch Utilization (%)
HLLD𝑖 () 29.7 1.9 1156 566 2.04 24.7 0.18 75 29 120 32 15 1.03 0.64
HLLD 𝑗,𝑘 () 110.6 13.8 312 248 1.26 24.7 0.05 52 43 108 32 4 0.28 0.17
Rec() 25.7 4.8 3584 417 8.60 46.8 0.74 85 25 92 32 58 1.51 4.53
RHS() 35.0 6.6 68 447 0.15 51.6 0.06 69 37 88 32 3 0.57 0.31
CT_EMF() 32.4 12.1 33 717 0.04 51.8 0.02 24 5 88 32 0.3 0.22 0.02
US() 119.8 22.4 3 290 0.01 71.1 0.07 37 15 56 32 4 0.19 0.75
Most kernels sustain substantial HBM traffic (up to 717 GB/s) but remain far below the 1.6 TB/s MI250X peak. Low VALU utilization (< 5%) and IPC (< 0.1) for memory-intensive kernels confirm latency-limited performance rather than bandwidth saturation. Only Reconstruct() achieves higher efficiency (IPC 0.74, VALU 58%, AI 8.6). L1 hit rates vary substantially: Reconstruct() ≈ 85%, CT_EMF() and US() below 40%; L2 rates remain modest. Irregular/strided access patterns amplify memory latency, limiting cache hierarchy effectiveness. HLLD() exhibits 3.7× performance gap between unit-stride (𝑖-sweep: 29.7 ms) and strided sweeps (𝑗, 𝑘: 110.6 ms). This correlates with L1 hit rates (75% vs. 52%) and VALU utilization (15% vs. 4%). The array layout causes 𝑗, 𝑘 directions to bypass L1, so increased VMEM latency causes wavefronts stall. On A100, the same anisotropy is mild (approximately 20%), while on MI250X it dominates performance. This can be seen as a confirmation of its latency-bound behavior: despite substantial HBM traffic relative to the 1.6 TB/s per-GCD peak on MI250X, performance does not improve under strided patterns. RHS() shows 33% directional variation (vs. 14% on A100). Achieved occupancy spans 25–71%, yet higher occupancy does not improve performance: US() achieves 71% occupancy, while RHS() and CT_EMF() reach 50%, but all exhibit low IPC and VALU utilization. Occupancy alone cannot hide
4
Discussion
48.1 Runtime [s]
gPLUTO Performance on LUMI (AMD MI250X)
40 20
13.9
17.2
0 A100 ACC
A100 OMP
MI250X OMP
Figure 1: Total execution time for 10 steps of 3D Orszag-Tang test.
Other metrics correlate with this behavior. Occupancy alone cannot hide latency without coalescing: enabled by its larger CUlevel register file, low VMEM utilization and modest L2 hit rates indicate frequent wavefront stalls rather than instruction bottlenecks. Although IPC is not directly comparable across architectures, it provides a qualitative indicator that confirms this behavior. High branch efficiency on both platforms rules out divergence as a cause, confirming that memory coalescing and latency hiding drive the MI250X slowdown. Table 6 summarizes how C++ wrapper abstractions (Ary1D/Ary2D) affect performance across regimes. In 2D kernels, OpenMP’s conservative pointer handling triggers massive spilling and suboptimal grid mapping (47× slowdown). In 3D, 12 Kernel-level profiling (Tables 2 and 5) reports single-invocation timings under iso-
lated conditions, while Figure 1 aggregates several kernel launches per time step across spatial directions, RK stages, and auxiliary routines. Absolute times are not directly comparable.
On the Limits of Performance Portability in Directive-Based GPU Programming
higher parallelism hides latency, recovering near-parity (6% gap) with OpenACC through the unified CUDA backend. On MI250X, the larger CU-level register file mitigates spilling, shifting the bottleneck to memory coalescing, further amplified by the relative immaturity of OpenMP offloading on ROCm (AOCC/Clang) backend for allocation/alias analysis. Recovering A100-level performance therefore requires loop reordering, data-layout transformations, and direction-specific specialization for register management. Such interventions are well beyond straightforward pragma translation. Table 6: C++ abstraction (Ary1D/Ary2D) impact on kernel behavior across architectures and programming models. Configuration A100 OpenACC 2D A100 OpenMP 2D A100 OpenACC 3D A100 OpenMP 3D MI250X OpenMP 3D
Register Usage 104 reg/thread 98 reg/thread 172 reg/thread 182 reg/thread 108-120 VGPR/wf
Spilling None Massive Moderate Moderate Minimal
Primary Bottleneck Memory latency Spilling + grid Latency + spilling Latency (hidden) Strided access
Performance Fast 47 × slower Moderate 6% slower 3 × slower
These effects are particularly visible in smaller kernels on A100 as well (Section 3.1.1), where OpenMP overheads (mapping, team setup, privatization) combine with complex C++ wrappers to trigger excessive register pressure and spilling. In the 3D Orszag–Tang problem, higher parallelism hides spilling latency, achieving nearparity despite moderate register pressure in individual kernels. Notably, the observed occupancy range reflects architectural and compiler differences rather than insufficient parallelism. An alternative approach to mitigate register pressure would involve explicitly constraining register usage through compiler flags (e.g., -maxrregcount on NVHPC) to increase occupancy at the cost of additional spilling. However, optimal limits are architecture- and kernel-specific, undermining portability. We therefore prefer algorithmic transformations that reduce intrinsic register pressure. Finally, an additional factor is unified memory on MI250X, whose automatic page migration may amplify strided-access overhead. Isolating this contribution for a pointer-intensive C++ code like gPLUTO requires explicit allocation strategies that underscore practical portability challenges.
5
Conclusions
This work assessed performance portability of directive-based GPU programming by porting the production astrophysical code gPLUTO to NVIDIA A100 and AMD MI250X GPUs using OpenACC and OpenMP. Directive-based models ease incremental porting of legacy codes, but require architecture-aware tuning, detailed profiling, and data layout transformations to achieve competitive performance. A key outcome is that performance portability is governed more by compiler maturity and backend behavior than by the programming model alone. The NVHPC toolchain benefits from over a decade of CUDA-targeted optimization heuristics, while the ROCm OpenMP implementation, though based on the same LLVM infrastructure, represents a more recent engineering effort with GPUspecific passes still under active development. On NVIDIA A100, the shared CUDA backend enables near-parity between OpenACC and OpenMP. In contrast, the same OpenMP implementation and workload on AMD MI250X exhibits a slowdown of ∼ 3×. These are not explained by insufficient parallelism, but arise from strong sensitivity to strided memory-access patterns, limited latency hiding,
Conference’17, July 2017, Washington, DC, USA
and variable compiler optimization across kernels. We observe that architectural bottlenecks manifest differently across vendors: on MI250X, high sustained HBM bandwidth does not translate into throughput for irregular or directionally dependent kernels, while on NVIDIA GPUs C++ abstractions primarily manifest through register pressure and spilling. We also show that application-level performance is dictated by a small number of latency-sensitive kernels: large speedups in secondary kernels (e.g., Reconstruct()) do not compensate for slowdowns along the critical path. Overall, porting from OpenACC on NVIDIA to OpenMP on AMD is not a "search and replace" process: recovering performance requires intrusive transformations such as loop reorderings, datalayout changes, software transpositions, and direction-specific specializations. Future work should also isolate the impact of unified versus managed memory on AMD platforms and explore explicit data-movement strategies that avoid page-migration overhead while preserving maintainability. Furthermore, extending the analysis to Intel GPUs would provide a third vendor perspective, though system-dependent unified_shared_memory (USM) behavior and compiler maturity introduce non-trivial engineering challenges [14] (mappers, explicit allocation APIs, compiler-specific paths). Moreover, inspecting OpenMP implementations with alternative compilers (e.g., upstream Clang/LLVM) on NVIDIA hardware to isolate the impact of NVHPC-specific optimizations from intrinsic OpenMP model characteristics would be a valuable insight. Quantifying porting effort alongside performance gain, e.g., refactoring complexity per kernel, would further make these findings more actionable for the HPC community. While such work increases development cost, multi-vendor support remains essential for broad access to EuroHPC systems, long-term sustainability of scientific codes, and exploitation of architecture-specific strengths. Combining metricbased approaches with detailed profiling therefore represents an important direction for future research. In summary, correctness portability across modern GPUs is readily achievable, whereas performance portability remains a significant challenge.
Acknowledgments All authors contributed equally. The authors acknowledge M. Bettencourt and G. Rossi for their collaboration. This work was supported by the SPACE Centre of Excellence (EU Grant No. 101093441). We acknowledge ISCRA for awarding access to the LEONARDO supercomputer, owned by the EuroHPC Joint Undertaking, hosted by CINECA (Italy); and the EuroHPC Joint Undertaking for access to LUMI, hosted by CSC (Finland) and the LUMI consortium. AI assistants (ChatGPT, Claude) were used solely for light editing and proofreading; all scientific content is the original work of the authors.
References [1] Advanced Micro Devices, Inc. 2023. Omniperf: Performance Analysis Tool for AMD GPUs. https://github.com/ROCm/rocm-systems. [2] Advanced Micro Devices, Inc. 2024. GPU Architecture Hardware Specifications (ROCm Documentation). https://rocm.docs.amd.com/en/docs-6.0.2/reference/ gpu-arch/gpu-arch-spec-overview.html [3] M. Aldinucci et al. 2021. Practical parallelization of scientific applications with OpenMP, OpenACC and MPI. J. Parallel and Distrib. Comput. 157 (2021), 13–29. doi:10.1016/j.jpdc.2021.05.017
Conference’17, July 2017, Washington, DC, USA
[4] S. F. Antao et al. 2016. Offloading Support for OpenMP in Clang and LLVM. In 2016 Third Workshop on the LLVM Compiler Infrastructure in HPC (LLVM-HPC). 1– 11. doi:10.1109/LLVM-HPC.2016.006 [5] Argonne Leadership Computing Facility. 2021. Inside the NVIDIA Ampere A100 GPU. Slide deck. https://www.alcf.anl.gov/sites/default/files/2021-07/ALCF_ A100_20210728%5B80%5D.pdf [6] Bertolli, C. et al. 2015. Integrating GPU support for OpenMP offloading directives into clang. In Proceedings of LLVM-HPC 2015. Association for Computing Machinery, Inc. doi:10.1145/2833157.2833161 [7] J. Choquette et al. 2021. NVIDIA A100 GPU: Performance and Innovation. IEEE Micro 41, 2 (2021), 29–35. doi:10.1109/MM.2021.3061394 [8] J. H. Davis et al. 2025. Taking GPU Programming Models to Task for Performance Portability. In Proceedings of the 39th ACM International Conference on Supercomputing (ICS ’25). ACM, 776–791. doi:10.1145/3721145.3730423 [9] T. Deakin et al. 2020. Performance Portability across Diverse Computer Architectures. In 2019 IEEE/ACM International Workshop on Performance, Portability and Productivity in HPC (P3HPC). Institute of Electrical and Electronics Engineers (IEEE). doi:10.1109/P3HPC49587.2019.00006 [10] T. Deakin and T. G. Mattson. 2023. Programming Your GPU with OpenMP: Performance Portability for GPUs. MIT Press. https://mitpress.mit.edu/ 9780262547536/programming-your-gpu-with-openmp/ [11] T. Deakin, J. Price, M. Martineau, and S. McIntosh-Smith. 2018. Evaluating attainable memory bandwidth of parallel programming models via BabelStream. International Journal of Computational Science and Engineering 17, 3 (2018), 247–262. doi:10.1504/IJCSE.2018.095847 [12] A. Dubey et al. 2021. Performance Portability in the Exascale Computing Project: Exploration Through a Panel Series. Computing in Science & Engineering 23, 5 (2021), 46–54. doi:10.1109/MCSE.2021.3098231 [13] H. C. Edwards and C. R. Trott. 2013. Kokkos: Enabling Performance Portability Across Manycore Architectures. In 2013 Extreme Scaling Workshop (xsw 2013). 18–24. doi:10.1109/XSW.2013.7 [14] W. Elwasif. 2023. Experimental Characterization of OpenMP Offloading Memory Operations and Unified Shared Memory Support. In OpenMP: Advanced Task-Based, Device and Compiler Programming. Springer Nature Switzerland, Cham, 210–225. doi:10.1007/978-3-031-40744-4_14 [15] ENCCS. 2022. Hierarchical Roofline Performance Analysis on AMD GPUs. https: //enccs.github.io/amd-rocm-development. [16] A. Folch et al. 2023. The EU Center of Excellence for Exascale in Solid Earth (ChEESE): Implementation, results, and roadmap for the second phase. Future Generation Computer Systems 146 (2023), 47–61. doi:10.1016/j.future.2023.04.006 [17] Y. Fridman, Y. Goren, and G. Oren. 2025. From OpenACC to OpenMP5 GPU Offloading: Performance Evaluation on NAS Parallel Benchmarks. In Proceedings of the 2025 4th International Workshop on Extreme Heterogeneity Solutions (ExHET ’25). Association for Computing Machinery, New York, NY, USA, 10–18. doi:10.1145/3720555.3721989 [18] A. Garcia et al. 2025. MaX - Materials Design at the eXascale: Recent Selected Results. In Proceedings of the 22nd ACM International Conference on Computing Frontiers (CF ’25). 150–156. doi:10.1145/3706594.3727577 [19] P. Grete, F. W. Glines, and B. W. O’Shea. 2021. K-Athena: A Performance Portable Structured Grid Finite Volume Magnetohydrodynamics Code. IEEE Transactions on Parallel and Distributed Systems 32, 1 (2021), 85–97. doi:10.1109/TPDS.2020. 3010016 [20] M. A. Heroux and J. M. Willenbring. 2009. Barely sufficient software engineering: 10 practices to improve your CSE software. In 2009 ICSE Workshop on Software Engineering for Computational Science and Engineering. 15–21. doi:10.1109/ SECSE.2009.5069157 [21] J. K. Holmen, B. Peterson, and M. Berzins. 2019. An Approach for Indirectly Adopting a Performance Portability Layer in Large Legacy Codes. In 2019 IEEE/ACM International Workshop on Performance, Portability and Productivity in HPC (P3HPC). 36–49. doi:10.1109/P3HPC49587.2019.00009 [22] M. Khalilov and A. Timoveev. 2021. Performance analysis of CUDA, OpenACC and OpenMP programming models on TESLA V100 GPU. Journal of Physics: Conference Series 1740, 1 (jan 2021), 012056. doi:10.1088/1742-6596/1740/1/ 012056 [23] M. Klemm. 2025. OpenMP® Target Offloading for AMD Instinct GPUs and APUs. https://tu-dresden.de/zih/das-department/ressourcen/dateien/ kolloquium/2025_03_27-MichaelKlemm.pdf. Tutorial on OpenMP offloading and GPU performance, Accessed 2025. [24] E. Krishnasamy et al. 2026. Performance and Programmability of MPI+X Integration with CUDA, HIP, SYCL, OpenACC, and OpenMP Offloading for Supercomputing: A Case Study on Dense Matrix-Vector Multiplication. doi:10.1145/ 3784828.3786264 [25] A. Marowka. 2025. Portability efficiency approach for calculating performance portability. Future Generation Computer Systems 170 (2025), 107826. doi:10. 1016/j.future.2025.107826 [26] N. A. Mehta, R. Gayatri, Y. Ghadar, C. Knight, and J. Deslippe. 2021. Evaluating Performance Portability of OpenMP for SNAP on NVIDIA, Intel, and AMD GPUs Using the Roofline Methodology. In Accelerator Programming Using Directives.
Alessandro Romeo, Nitin Shukla, Stefano Truzzi, Alessio Suriano, and Andrea Mignone
Springer International Publishing, Cham, 3–24. doi:10.1007/978-3-030-74224-9_1 [27] S. Memeti, L. Li, S. Pllana, J. Kołodziej, and C. Kessler. 2017. Benchmarking OpenCL, OpenACC, OpenMP, and CUDA: Programming Productivity, Performance, and Energy Consumption. In Proceedings of the 2017 Workshop on Adaptive Resource Management and Scheduling for Cloud Computing (Washington, DC, USA) (ARMS-CC ’17). Association for Computing Machinery, New York, NY, USA, 1–6. doi:10.1145/3110355.3110356 [28] A. Myers et al. 2021. Porting WarpX to GPU-accelerated platforms. Parallel Comput. 108 (2021), 102833. doi:10.1016/j.parco.2021.102833 [29] NVIDIA. 2026. NVIDIA Ampere GPU Architecture Tuning Guide. https://docs. nvidia.com/cuda/ampere-tuning-guide/index.html [30] NVIDIA Corporation. 2023. Nsight Compute Documentation: Memory Workload Analysis. https://docs.nvidia.com/nsight-compute/NsightCompute/index.html. [31] OpenACC-Standard.org. 2023. The OpenACC Application Programming Interface, Version 3.3. Technical Report. OpenACC Organization. https: //www.openacc.org/specification [32] OpenMP Architecture Review Board. 2021. OpenMP Application Programming Interface, Version 5.2. Technical Report. OpenMP ARB. https://www.openmp. org/wp-content/uploads/OpenMP-API-Specification-5-2.pdf [33] J. Owens et al. 2008. GPU computing. Proc. IEEE 96 (05 2008), 879–899. doi:10. 1109/JPROC.2008.917757 [34] S. J. Pennycook, J. D. Sewall, and V. W. Lee. 2016. A Metric for Performance Portability. arXiv:1611.07409 [cs.PF] https://arxiv.org/abs/1611.07409 [35] M. Rossazza et al. 2026. The PLUTO code on GPUs: A first look at Eulerian MHD methods. Astronomy and Computing (2026), 101076. doi:10.1016/j.ascom.2026. 101076 [36] G. Schieffer et al. 2024. Understanding Data Movement in AMD Multi-GPU Systems with Infinity Fabric. In Proceedings of the SC ’24 Workshops of the International Conference on High Performance Computing, Network, Storage, and Analysis (Atlanta, GA, USA) (SC-W ’24). IEEE Press, 567–576. doi:10.1109/ SCW63240.2024.00079 [37] J. Sewall, S. J. Pennycook, D. Jacobsen, T. Deakin, and S. McIntosh-Smith. 2020. Interpreting and Visualizing Performance Portability Metrics. In 2020 IEEE/ACM International Workshop on Performance, Portability and Productivity in HPC (P3HPC). 14–24. doi:10.1109/P3HPC51967.2020.00007 [38] N. Shukla et al. 2025. Towards Exascale Computing for Astrophysical Simulation Leveraging the Leonardo EuroHPC System. Procedia Computer Science 267 (2025), 112–123. doi:10.1016/j.procs.2025.08.238 Proceedings of the Third EuroHPC user day. [39] N. Shukla et al. 2026. Exascale computing to accelerate discoveries in astrophysics and space plasma physics. Nature Astronomy 10 (2026), 330–334. doi:10.1038/ s41550-026-02807-8 [40] C. P. Sishtla et al. 2019. Multi-GPU Acceleration of the iPIC3D Implicit Particlein-Cell Code. In Computational Science – ICCS 2019. Springer International Publishing, Cham, 612–618. [41] A. Smith and N. James. 2022. AMD Instinct™ MI200 Series Accelerator and Node Architectures. In 2022 IEEE Hot Chips 34 Symposium (HCS). 1–23. doi:10.1109/ HCS55958.2022.9895477 [42] J. M. Stone, K. Tomida, C. J. White, and K. G. Felker. 2020. The Athena++ Adaptive Mesh Refinement Framework: Design and Magnetohydrodynamic Solvers. The Astrophysical Journal Supplement Series 249, 1 (June 2020), 4. doi:10.3847/15384365/ab929b [43] A. Suriano et al. 2026. The PLUTO code on GPUs: Offloading Lagrangian Particle methods. Astronomy and Computing 55 (2026), 101088. doi:10.1016/j.ascom. 2026.101088 [44] S. Tandon et al. 2024. Porting HPC Applications to AMD Instinct™ MI300A using Unified Memory and OpenMP®. In ISC High Performance 2024 Research Paper Proceedings (39th International Conference). 1–9. doi:10.23919/ISC.2024. 10528925 [45] S. Wienke, P. Springer, C. Terboven, and D. Mey. 2012. OpenACC - First Experiences with Real-World Applications. In Euro-Par 2012 Parallel Processing. Springer Berlin Heidelberg, Berlin, Heidelberg, 859–870. [46] S. Wienke, C. Terboven, J. C. Beyer, and M. S. Müller. 2014. A Pattern-Based Comparison of OpenACC and OpenMP for Accelerator Computing. In Euro-Par 2014 Parallel Processing. Springer International Publishing, Cham, 812–823. [47] J. Williams et al. 2024. Optimizing BIT1, a Particle-in-Cell Monte Carlo Code, with OpenMP/OpenACC and GPU Acceleration. Springer Nature Switzerland, Cham, 316–330. doi:10.1007/978-3-031-63749-0_22 [48] S. Williams, A. Waterman, and D. Patterson. 2009. Roofline: An Insightful Visual Performance Model for Multicore Architectures. Commun. ACM 52, 4 (2009), 65–76. doi:10.1145/1498765.1498785 [49] Y. Yan et al. 2025. OpenMP: Balancing Productivity and Performance Portability. Springer. doi:10.1007/978-3-032-06343-4