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, which includes hipcc and related commands;
linear algebra libraries such as rocblas, rocsparse, hipblas, and hipsparse;
the rocHPL benchmark;
a ROCm-enabled variant of Open MPI;
tools such as rocprofiler and roctracer.

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-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.
We create two MPI processes on the node (mpirun -n 2). The --mca pml ucx flag ensures Open MPI selects ucx 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:

Heap map showing the bandwidth between every pair of GCDs.

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-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.
-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 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.

HIP and ROCm as first-class citizens