Building GCC with support for NVIDIA PTX offloading

GCC can offload C, C++, and Fortran code to an accelerator when using OpenACC or OpenMP where the code to offload is controlled by adding #pragma statements (or magic comments for Fortran), such as

#pragma acc kernels
for (int j = 1; j < n-1; j++) {
  for (int i = 1; i < m-1; i++) {
    Anew[j][i] = 0.25f * (A[j][i+1] + A[j][i-1] + A[j-1][i] + A[j+1][i]);
    error = fmaxf(error, fabsf(Anew[j][i] - A[j][i]));
  }
}

This blog post describes what I needed to do in order to build a GCC 7.2 compiler with support for offloading to NVIDIA GPUs on Ubuntu 16.10.

The first step is to install the NVIDIA CUDA toolkit. Googling shows lots of strange suggestions about what you need to do in order to get this to work (blacklisting drivers, adding the PCI address of your video card to config files, etc.), but it worked fine for me to just download the “deb (local)” file, and install it as

sudo dpkg -i cuda-repo-ubuntu1604-8-0-local-ga2_8.0.61-1_amd64.deb
sudo apt-get update
sudo apt-get install cuda

The toolkit is installed in /usr/local/cuda, and /usr/local/cuda/bin must be added to PATH so that GCC may find the ptxas tool.

The script below fetches the source code and builds the compiler and tools

Add $install_dir/lib64 to LD_LIBRARY_PATH, and the compiler can now be used to offload OpenACC code by compiling as

$install_dir/bin/gcc -O3 -fopenacc test.c

or OpenMP as

$install_dir/bin/gcc -O3 -fopenmp test.c

You may need to pass -foffload=-lm to the compiler if the code you offload contains math functions that cannot be directly generated as PTX instructions.

Updated 2017-12-23: Changed the script to build GCC 7.2 instead of trunk as there are some problems with the trunk compiler at the moment...