“Fatbinaries” for Mixed Clusters

Mixed clusters are machines where different nodes provide different compute capabilities. At the moment these kinds of clusters are common in heterogeneous computing, where some nodes only provide multi-core CPUs and some nodes additionally provide GPUs as well.

In building software for such machines one could build one version just for CPU nodes, and another version for the GPU nodes. The downsides of doing this include:

  • you need to build the code twice

  • each build needs separate compiler flags to target the right architecture and the build procedure needs to be adapted accordingly

  • afterwards to need to make sure you run the right executable on the right node

It would be much easier if you could compile a program only once so that it includes code for both CPUs and GPUs, and the program selects the code it executes based on what hardware it finds on a given node. This is not a new idea and a number of ways of implementing this have been developed. One way is to generate proper fatbinaries. These programs include multiple versions for different hardware in a single binary. Another way to provide this is to compile the code to an intermediate representation, and use a just-in-time compiler to generate the binary for the hardware present at runtime. Either way in both cases you have a single executable that can run on multiple hardware architectures. Here all such approaches are referred to as fatbinaries.

Prerequisites for generating fatbinaries

In order to pursue a fatbinary approach the code has to have parts that can be adapted to different hardware. As GPUs are highly parallel processing units a code needs to identify parts that are suitable for parallelization. In addition those parts need to be such that they can be adapted to different kinds of hardware. These prerequisites make codes that are parallelised with OpenMP or OpenACC (possibly in addition to MPI) most suitable.

OpenMP is a directive based approach where special comments in the code provide hints to the compiler about what parts are suitable for parallelisation. OpenMP is developed to parallelise codes in a multithreaded way, and it was created long before GPUs became important. Therefore, with OpenMP it is important to check whether the program uses off-loading as this is the mechanism to place compute on a GPU. Without off-loading an OpenMP program is a CPU only code.

OpenACC is also a directive based approach but it has been specifically developed to target GPU based systems. You can build OpenACC codes to run on CPUs as well but the distribution of compute across cores is not always what you would expect. Therefore, we will focus on OpenMP codes below.

Compiler suites

On amplitUDE we have installed a particular suite of compilers that can generate fatbinaries. This is a special build of GCC which is available as the gcc/15.2.0 module. Actually, this build of GCC uses Nvidia’s ptxas under the hood to generate GPU code, the intermediate representation of your program that feeds into that is just generated with the GNU compilers.

An example

Let’s look at a simple example, build and run it, and look at the profile data to see what resources it is using. We’ll look at both compiler suites in turn.

As a simple example consider a simple matrix-matrix multiplication. This is a common operation in HPC codes that is well suited to parallelisation. These small codes are specifically written so you can easily copy-and-paste them to try it out for yourself.

The C-code lives in mxm_c_openmp.c and contains:

#include <stdio.h>
#include <stdlib.h>
#include <omp.h>

int main(int argc, char *argv[]) {
    #define nsize 8000
    static double a[nsize][nsize];
    static double b[nsize][nsize];
    static double c[nsize][nsize];
    size_t num_per_thread;
    size_t num_devices;
    size_t max_threads;
    size_t num_procs;
    size_t kdev; // The device identifier
    size_t ii, jj, kk; // loop counters
    size_t iilo, iihi, iilen;
    num_devices = omp_get_num_devices();
    max_threads = omp_get_max_threads();
    if (num_devices > 0) {
        num_procs = num_devices;
    }
    else {
        num_procs = max_threads;
    }
    if (num_procs == 0) {
        printf("ERROR: too few threads for the number of devices");
        exit(200);
    };
    printf("mxm_c_openmp.c: num_procs = %d\n",num_procs);
    // Initialize the matrices
    for (ii = 0; ii < nsize; ii++) {
        for (jj = 0; jj < nsize; jj++) {
             a[ii][jj] = ((double)1.0)/((double)(jj*nsize+ii));
             b[ii][jj] = ((double)1.0)/((double)(ii*nsize+jj));
        }
    }
    // Now do the compute
    num_per_thread = (nsize+num_procs-1)/num_procs;
    #pragma omp parallel for private(iilo, iihi, iilen)
    for (kdev = 0; kdev < num_procs; kdev++) {
        iilo = kdev*num_per_thread;
        iihi = (kdev+1)*num_per_thread <= nsize ? (kdev+1)*num_per_thread : nsize;
        iilen = iihi - iilo;
        #pragma omp target device(kdev) map(from: c[iilo:iilen][:]) map(to: a[iilo:iilen][:], b, iilen) nowait
        {
            #pragma omp teams distribute parallel for collapse(2) default(shared) firstprivate(iilen) private(ii, jj, kk)
            for (ii = 0; ii < iilen; ii++) {
                 for (jj = 0; jj < nsize; jj++) {
                      c[ii][jj] = 0.0;
                      for (kk = 0; kk < nsize; kk++) {
                          c[ii][jj] += a[ii][kk] * b[kk][jj];
                      }
                 }
             }
        }
    }
    #pragma omp taskwait
    printf("mxm_c_openmp.c: finished\n");
}

Alternatively Fortran is a commonly used language in the HPC domain. The same algorithm as above is kept in mxm_f_openmp.F90 and reads:

program mxm_f_openmp
  use omp_lib
  ! Compute C = A*B
  implicit none
  integer, parameter :: dp = selected_real_kind(15,3)
  integer, parameter :: nsize = 8000
  real(kind=dp) :: a(nsize,nsize)
  real(kind=dp) :: b(nsize,nsize)
  real(kind=dp) :: c(nsize,nsize)
  integer :: num_per_thread
  integer :: max_threads
  integer :: num_devices
  integer :: num_procs
  integer :: kdev ! The device identifier
  integer :: ii, jj, kk ! The loop counters
  integer :: jjlo, jjhi
  num_devices = omp_get_num_devices()
  max_threads = omp_get_max_threads()
  if (num_devices.gt.0) then
    num_procs = num_devices
  else
    num_procs = max_threads
  endif
  if (num_procs.le.0) then
    write(*,*)" ERROR: too few resources for the calculation"
    stop 200
  endif
  write(*,*)"mxm_f_openmp.F90: num_procs = ",num_procs
  ! Initialize matrices
  do jj = 1, nsize
    do ii = 1, nsize
      a(ii,jj) = 1.0_dp/(jj*nsize+ii)
      b(ii,jj) = 1.0_dp/(ii*nsize+jj)
    enddo
  enddo
  ! now do the compute
  num_per_thread = (nsize+num_procs-1)/num_procs
  !$omp parallel do private(jjlo, jjhi)
  do kdev = 0, num_procs-1
    jjlo = kdev*num_per_thread+1
    jjhi = min((kdev+1)*num_per_thread,nsize)
    !$omp target device(kdev) map(from: c(:,jjlo:jjhi)) map(to: a, b(:,jjlo:jjhi), jjlo, jjhi) nowait
    !$omp teams distribute parallel do collapse(2) default(shared) firstprivate(jjlo, jjhi) private(jj,ii,kk)
    do jj = jjlo, jjhi
      do ii = 1, nsize
        c(ii,jj) = 0.0_dp
        do kk = 1, nsize
          c(ii,jj) = c(ii,jj) + a(ii,kk) * b(kk,jj)
        enddo
      enddo
    enddo
    !$omp end teams distribute parallel do
    !$omp end target
  enddo
  !$omp end parallel do
  !$omp taskwait
  write(*,*)"mxm_f_openmp.F90: finished"
end program mxm_f_openmp

Using the GCC suite

To compile the code run the following bash script:

#!/bin/bash
module load nvhpc/25.11-cuda13
module load gcc/15.2.0
#
gcc      -fopenmp  mxm_c_openmp.c    -o mxm_c_openmp_gcc
#
gfortran -fopenmp  mxm_f_openmp.F90  -o mxm_f_openmp_gcc

With the GCC compilers we just need to provide the -fopenmp flag to translate the OpenMP code.

To run the executables on a CPU node submit the following batch script:

#!/bin/bash
#SBATCH -t 1:00:00
#SBATCH -N 1
#SBATCH --tasks-per-node=1
#SBATCH --cpus-per-task=112

module load nvhpc/25.11-cuda13
module load gcc/15.2.0

nsys profile --output mxm_f_openmp_gcc_cpu.nsys-rep  ./mxm_f_openmp_gcc
nsys profile --output mxm_c_openmp_gcc_cpu.nsys-rep  ./mxm_c_openmp_gcc

The nsys command invokes the Nvidia Nsight System profiling tool. It records the behavior of a program and stores that data in a Nsys report (a nsys-rep file). The data in these files can be visualized with Nvidia Nsight Systems visualizer. You need to run the visualizer on your local machine and you need to have Nvidia’s Nsight System installed.

Likewise to the code on a GPU node submit the following batch script:

#!/bin/bash
#SBATCH -t 1:00:00
#SBATCH -N 1
#SBATCH --tasks-per-node=1
#SBATCH --cpus-per-task=4
#SBATCH --gres=gpu:4
#SBATCH --partition=GPU-big

export OMP_TARGET_OFFLOAD=MANDATORY

module load nvhpc/25.11-cuda13
module load gcc/15.2.0

nsys profile --output mxm_f_openmp_gcc_gpu.nsys-rep  ./mxm_f_openmp_gcc
nsys profile --output mxm_c_openmp_gcc_gpu.nsys-rep  ./mxm_c_openmp_gcc

Subsequently we can vizualize the profiling data with the Nvidia Nsight Systems visualizer. For the jobs run on CPU nodes we get

the performance profile of the Fortran matrix-matrix multiplication with OpenMP compiled with GCC on CPUs the performance profile of the C matrix-matrix multiplication with OpenMP compiled with GCC on CPUs

The figures show that on CPU nodes the code is executed on multiple cores.

Likewise we can visualize the behavior of the code on GPU node. For those calculations we get:

the performance profile of the Fortran matrix-matrix multiplication with OpenMP compiled with GCC on GPUs the performance profile of the C matrix-matrix multiplication with OpenMP compiled with GCC on GPUs

As one can see, on GPU nodes the compute is offloaded to the GPU devices. Note that we launched the same executables on the CPU as well as the GPU nodes.