--- title: "HIP and ROCm as first-class citizens" slug: hip-and-rocm-as-first-class-citizens url: https://listedarticles.com/articles/hip-and-rocm-as-first-class-citizens canonical_url: https://hpc.guix.info/blog/2026/10/hip-and-rocm-as-first-class-citizens/ content_type: blog_post language: en published_at: 2026-10-09T00:00:00.000Z updated_at: 2026-10-09T14:41:14.817Z author: "Ludovic Courtès" authored_by: human publisher: "Guix-HPC" publisher_url: https://hpc.guix.info/ topics: ["Open Source", "Linux", "Hardware", "Performance"] license: CC-BY-SA-4.0 word_count: 1485 reading_minutes: 6 citation: "Ludovic Courtès, Guix-HPC. \"HIP and ROCm as first-class citizens.\" 9 Oct 2026. https://hpc.guix.info/blog/2026/10/hip-and-rocm-as-first-class-citizens/ (CC-BY-SA-4.0)" # The full text follows. The web page shows an extract and sends readers # to the source above; quote the citation and link the canonical URL. --- # HIP and ROCm as first-class citizens > Ludovic Courtès reports that Guix developers migrated the whole AMD HIP/ROCm stack from the Guix-HPC channels into Guix proper starting with 7.1.1: more than 40 packages covering the toolchain, linear algebra libraries, rocHPL, a ROCm Open MPI and profiling tools, plus HIP variants of scientific software. # HIP and ROCm as first-class citizens Earlier this year, [two years after AMD engineers contributed packages the HIP/ROCm stack to the Guix-HPC channels](https://hpc.guix.info/blog/2024/01/hip-and-rocm-come-to-guix/), Guix developers migrated the whole HIP/ROCm stack into Guix proper, starting with version 7.1.1. This was quite a milestone as it makes HIP/ROCm first-class citizens and gives them more exposure and better support in the community. The package set consists of more than 40 packages covering many things: - [the toolchain itself](https://hpc.guix.info/package/rocm-toolchain) , which includes`hipcc` and related commands; - linear algebra libraries such as [`rocblas`](https://hpc.guix.info/package/rocblas) ,[`rocsparse`](https://hpc.guix.info/package/rocsparse) ,[`hipblas`](https://hpc.guix.info/package/hipblas) , and[`hipsparse`](https://hpc.guix.info/package/hipsparse) ; - the [rocHPL](https://hpc.guix.info/package/rochpl) benchmark; - a [ROCm-enabled variant of Open MPI](https://hpc.guix.info/package/openmpi-rocm) ; - tools such as [`rocprofiler`](https://hpc.guix.info/package/rocprofiler) and[`roctracer`](https://hpc.guix.info/package/roctracer) . We are also gradually adding HIP/ROCm variants of scientific software such as [CP2K](https://hpc.guix.info/package/cp2k-hip-rocm) and [Chameleon](https://hpc.guix.info/package/chameleon-hip-rocm), a dense linear algebra solver developed at Inria. # Selecting target GPUs The set of [AMD GPU architectures grows quickly](https://llvm.org/docs/AMDGPUUsage.html#amdgpu-processor-table). As packagers, we choose a default set of target GPU architectures to build ROCm/HIP-enabled applications for, but that set of architectures must be limited given the build time and size of resulting application binaries. It is crucial for users to be able to override this default set of target architectures to build specifically for the architecture(s) they want. To address that, we added a new package transformation option to Guix [called `--amd-gpu`](https://guix.gnu.org/manual/devel/en/html_node/Package-Transformation-Options.html#index-AMD-GPUs). Just like [`--tune` lets you build a package optimized for a specific CPU micro-architecture](https://hpc.guix.info/blog/2022/01/tuning-packages-for-a-cpu-micro-architecture/), `--amd-gpu` creates, on the fly, a variant of the relevant packages built specifically for the given GPU architecture(s). For example, here is how you would run a variant of the [ROCm bandwidth test](https://hpc.guix.info/package/rocm-bandwidth-test) built specifically for AMD Instinct MI250 (`gfx90a`) and for AMD Instinct MI300 (`gfx942`): ``` $ srun --tasks-per-node=1 -N1 --exclusive … \ guix shell rocm-bandwidth-test --amd-gpu=gfx90a,gfx942 -- \ rocm-bandwidth-test plugin --run tb p2p TransferBench v1.64.00 … Bytes Per Direction 268435456 Unidirectional copy peak bandwidth GB/s [Local read / Remote write] (GPU-Executor: GFX) SRC+EXE\DST CPU 00 CPU 01 CPU 02 CPU 03 GPU 00 GPU 01 GPU 02 GPU 03 CPU 00 -> 18.85 18.55 18.36 18.45 19.00 18.61 18.69 18.70 CPU 01 -> 7.87 8.82 8.42 8.63 7.99 8.16 7.40 7.64 CPU 02 -> 6.30 6.53 6.42 7.05 5.58 5.93 5.94 5.91 CPU 03 -> 6.34 7.09 7.50 7.35 6.17 6.23 6.37 6.23 GPU 00 -> 1307.15 90.22 90.85 91.45 1360.70 91.24 91.39 92.14 GPU 01 -> 90.58 1371.25 92.78 90.63 91.08 1431.93 92.38 90.84 GPU 02 -> 91.12 92.57 1355.95 91.82 90.68 92.22 1407.58 91.52 GPU 03 -> 91.41 91.08 91.66 1348.37 91.99 90.92 91.88 1468.65 CPU->CPU CPU->GPU GPU->CPU GPU->GPU Averages (During UniDir): 10.09 9.66 404.93 91.52 … ``` (This particular run was on a node with MI300 GPUs.) Of course these GPU architecture identifiers are, well, hard to grasp. You can find the full list [in the LLVM documentation](https://llvm.org/docs/AMDGPUUsage.html#amdgpu-processor-table); should you make a typo or select an architecture that the toolchain at hand doesn’t support, Guix lets you know about it without going any further: ``` $ guix build rocm-bandwidth-test --amd-gpu=forgot-the-name gnu/packages/llvm.scm:2283:2: error: compiler rocm-toolchain@7.1.1 does not support AMD GPU target forgot-the-name hint: Compiler rocm-toolchain@7.1.1 supports the following AMD GPU targets: bonaire, carrizo, fiji, generic, generic-hsa, gfx10-1-generic, gfx10-3-generic, gfx1010, gfx1011, gfx1012, gfx1013, gfx1030, gfx1031, gfx1032, gfx1033, gfx1034, gfx1035, gfx1036, gfx11-generic, gfx1100, gfx1101, gfx1102, gfx1103, gfx1150, gfx1151, gfx1152, gfx1153, gfx12-generic, gfx1200, gfx1201, gfx600, gfx601, gfx602, gfx700, gfx701, gfx702, gfx703, gfx704, gfx705, gfx801, gfx802, gfx803, gfx805, gfx810, gfx9-4-generic, gfx9-generic, gfx900, gfx902, gfx904, gfx906, gfx908, gfx909, gfx90a, gfx90c, gfx940, gfx941, gfx942, gfx950, hainan, hawaii, iceland, kabini, kaveri, mullins, oland, pitcairn, polaris10, polaris11, stoney, tahiti, tonga, tongapro, verde ``` # Performance So far, microbencharks are giving us green lights. Let’s take a closer look at benchmarks we ran on nodes with 8 MI250X GPUs of the [Adastra supercomputer](https://dci.dci-gitlab.cines.fr/webextranet/architecture/index.html). ## ROCm support for Open MPI First, there’s the bandwidth measured for MPI transfers among Graphics Compute Dies (GCDs)—specifically, using the ROCm-enabled Open MPI package, [`openmpi-rocm`](https://hpc.guix.info/package/openmpi-rocm). For this, we run the [OSU Micro-Benchmarks](https://hpc.guix.info/package/osu-micro-benchmarks-rocm) linked against `openmpi-rocm`, asking it to measure device-to-device transfers; we do that with large messages (16 MiB) and for all GCD pairs within a node, where each node has 8 GCDs: ``` export HSA_ENABLE_SDMA=0 for first in $(seq 0 7) do for second in $(seq $(($first + 1)) 7) do export HIP_VISIBLE_DEVICES="$first,$second" echo "# HIP_VISIBLE_DEVICES: $HIP_VISIBLE_DEVICES" guix time-machine -q --commit=f5c2937cddd4c8427f15b8b711f8a82211a09407 -- \ shell openmpi-rocm osu-micro-benchmarks-rocm -- \ mpirun -n 2 --mca pml ucx \ osu_bw -m $((16*1024*1024)):$((16*1024*1024)) D D done done ``` Some explanations: - The `time-machine` part selects the commit, and thus*the entire software stack* , that we have tested—fewer moving pieces. The`shell` bit specifies the packages we need in our environment. The commit we selected here provides a stack with Open MPI 5.0.10 and HIP/ROCm 7.1.1. - Setting `HSA_ENABLE_SDMA=0` , which turns off use of System Direct Memory Access (SDMA) by the HIP/ROCm runtime,[gives higher throughput](https://gpuopen.com/learn/amd-lab-notes/amd-lab-notes-gpu-aware-mpi-readme/) . - We create two MPI processes on the node (`mpirun -n 2` ). The`--mca pml ucx` flag ensures Open MPI selects[ucx](https://hpc.guix.info/package/ucx-rocm) as its interconnect backend. - Last, we run `osu_bw` , the bandwidth benchmark, for device-to-device (`D D` ) transfers with messages of 16 MiB. Setting the`HIP_VISIBLE_DEVICES` right above allows us to ensure transfers are made between these two GCDs. This gives us the communication matrix below, showing the bandwidth for unidirectional copies from device to device: The bandwidth we observe between each pair of GCDs matches [the GCD topology and Infinity Fabric links](https://gpuopen.com/learn/amd-lab-notes/amd-lab-notes-gpu-aware-mpi-readme/#gpu-to-gpu-communication-options-); for example, peak bandwidth between GCD 0 and GCD 1 is roughly four times that between GCD 0 and GCD 2, and two times that between GCD 0 and GCD 6. ## Computing benchmark What about computing throughput? A good test is [rocHPL](https://hpc.guix.info/package/rochpl), the ROCm-enabled variant of the classical [high-performance LINPACK](). For double-precision (aka. “Binary64”) floating point operations, the theoretical peak performance is 23.93 TFlop/s; for nodes with 8 GCDs, we can thus expect at most 23.93 x 8 = 191.5 TFlop/s per node. To get as close as possible to peak performance, we must arrange to let rocHPL work on a matrix that occupies almost all the GCD memory, which is 512 GiB per node here. For a single 8-GCD node, a matrix of 256,000 rows and columns fills 95% of device memory, making it a good choice—in line with what the [rocHPL wiki](https://github.com/ROCm/rocHPL/wiki/Common-rocHPL-run-configurations) suggests. We can run it with an incantation along these lines: ``` COMMIT=bf5d83139d8d4d7aa2d737b1e13620941abff2a1 # Workaround until # is merged. export ROCM_SMI_LIB_PATH="$(readlink -f $(guix time-machine \ -q --commit=$COMMIT -- \ build rocm-smi-lib | grep -v -e -bin$)/lib/librocm_smi64.so)" guix time-machine -q --commit=$COMMIT -- \ shell rochpl openmpi-rocm --tune=znver3 -- \ sh -x mpirun_rochpl -P 2 -Q 4 -N 256000 --NB 512 ``` Explanations: - The `time-machine` bit once again allows us to pin Guix to the commit for which we’ve run this benchmark, while`shell` sets up the execution environment. For good measure, we use`--tune` to[tune CPU code for the micro-architecture we have at hand](https://hpc.guix.info/blog/2022/01/tuning-packages-for-a-cpu-micro-architecture/) . - `-P` and`-Q` specify the number of rows and columns of the MPI grid; the product corresponds to the number of GCDs on the node.`-N` specifies the matrix size, as discussed above. That gives us a throughput of 160 TFlop/s—below the theoretical peak, but to our knowledge comparable to what others observe on MI250X. # Future work This post gives an overview of where the HIP/ROCm stack is in Guix and how its performance can be validated. There are a number of things we are planning to do, starting with a [minor-version upgrade of the HIP/ROCm stack](https://codeberg.org/guix/guix/pulls/10257) before we dive into more recent versions. More importantly, we are working on consolidation the set of benchmarks we want to use to validate the stack. Ideally, we won’t limit ourselves to micro-benchmarks and instead look at scientific applications that are known to exercise more of the supercomputer capabilities—such as GROMACS or CP2K. Our goal would be able to run a set of benchmarks before any package upgrade in the ROCm and MPI stacks, drawing from what admins [at Inria](https://jdev26.sciencesconf.org/726473) and [at CINES](https://jdev26.sciencesconf.org/720413) have been doing for their own clusters. All this is a much broader endeavor. To be continued! # Acknowledgments These benchmarks and additional tests were performed on the Adastra supercomputer hosted by CINES. Many thanks to our colleagues at CINES and to Florent Pruvost at Inria for their help. Huge thanks to the people who contributed to the HIP/ROCm stack in Guix, in particular to David Elsing for handling the bulk of the migration from the Guix-HPC channel and for upgrading those packages. Unless otherwise stated, blog posts on this site are copyrighted by their respective authors and published under the terms of the [CC-BY-SA 4.0](https://creativecommons.org/licenses/by-sa/4.0/) license and those of the [GNU Free Documentation License](https://www.gnu.org/licenses/fdl-1.3.html) (version 1.3 or later, with no Invariant Sections, no Front-Cover Texts, and no Back-Cover Texts).