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, 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
hipccand related commands;- linear algebra libraries such as
rocblas,rocsparse,hipblas, andhipsparse; - the rocHPL benchmark;
a [ROCm-enabled variant of
Open MPI](https://hpc.guix.info/package/openmpi-rocm) ;
tools such as
rocprofilerandroctracer.
We are also gradually adding HIP/ROCm variants of scientific software such as CP2K and Chameleon, a dense linear algebra solver developed at Inria.
Selecting target GPUs
The set of AMD GPU architectures grows quickly. 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.
Just like --tune lets you build a package optimized for a specific
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 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; 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.
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.
For this, we run the OSU
Micro-Benchmarks
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-machinepart selects the commit, and thus*the entiresoftware stack* , that we have tested—fewer moving pieces. The
shellbit 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 DirectMemory Access (SDMA) by the HIP/ROCm runtime,gives higher throughput .
We create two MPI processes on the node (
mpirun -n 2). The--mca pml ucxflag ensures Open MPI selectsucx as its interconnectbackend.
Last, we run
osu_bw, the bandwidth benchmark, for device-to-device(
D D) transfers with messages of 16 MiB. Setting theHIP_VISIBLE_DEVICESright 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; 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, 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 suggests.
We can run it with an incantation along these lines:
COMMIT=bf5d83139d8d4d7aa2d737b1e13620941abff2a1
# Workaround until <https://codeberg.org/guix/guix/pulls/11429>
# 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-machinebit once again allows us to pin Guix to thecommit for which we’ve run this benchmark, while
shellsets up the execution environment. For good measure, we use--tunetotune CPU code for the micro-architecture we have at hand .-Pand-Qspecify the number of rows and columns of the MPIgrid; the product corresponds to the number of GCDs on the node.
-Nspecifies 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 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 and at CINES 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 license and those of the GNU Free Documentation License (version 1.3 or later, with no Invariant Sections, no Front-Cover Texts, and no Back-Cover Texts).