Build GCC with support for offloading to NVIDIA GPUs

build-gcc-offload-nvptx.sh

#!/bin/sh

#
# Build GCC with support for offloading to NVIDIA GPUs.
#

work_dir=$HOME/offload/wrk
install_dir=$HOME/offload/install

# Location of the installed CUDA toolkit
cuda=/usr/local/cuda

# Build assembler and linking tools
mkdir -p $work_dir
cd $work_dir
git clone https://github.com/MentorEmbedded/nvptx-tools
cd nvptx-tools
./configure \
    --with-cuda-driver-include=$cuda/include \
    --with-cuda-driver-lib=$cuda/lib64 \
    --prefix=$install_dir
make || exit 1
make install || exit 1
cd ..

# Set up the GCC source tree
git clone git://sourceware.org/git/newlib-cygwin.git nvptx-newlib
git clone --branch releases/gcc-11 git://gcc.gnu.org/git/gcc.git gcc
cd gcc
contrib/download_prerequisites
ln -s ../nvptx-newlib/newlib newlib
cd ..
target=$(gcc/config.guess)

# Build nvptx GCC
mkdir build-nvptx-gcc
cd build-nvptx-gcc
../gcc/configure \
    --target=nvptx-none --with-build-time-tools=$install_dir/nvptx-none/bin \
    --enable-as-accelerator-for=$target \
    --disable-sjlj-exceptions \
    --enable-newlib-io-long-long \
    --enable-languages="c,c++,fortran,lto" \
    --prefix=$install_dir
make -j`nproc` || exit 1
make install || exit 1
cd ..

# Build host GCC
mkdir build-host-gcc
cd  build-host-gcc
../gcc/configure \
    --enable-offload-targets=nvptx-none \
    --with-cuda-driver-include=$cuda/include \
    --with-cuda-driver-lib=$cuda/lib64 \
    --disable-bootstrap \
    --disable-multilib \
    --enable-languages="c,c++,fortran,lto" \
    --prefix=$install_dir
make -j`nproc` || exit 1
make install || exit 1
cd ..

Hello, first of all thank's for your work. It compile perfectly. I just have a question, when I tr to offload on my GPU, the program doesn't launch on my GPU and I don't understand why.

Off course I have
#pragma omp target map(to:a[0:n]) map(to:b[0:n]) map(tofrom:c[0:n]) #pragma omp parallel for simd
before my loop. Do you have any ideas where it can come from ? Thank's for your help !

Off course I have #pragma omp target map(to:a[0:n]) map(to:b[0:n]) map(tofrom:c[0:n]) #pragma omp parallel for simd before my loop. Do you have any ideas where it can come from ? Thank's for your help !

I have never played with OpenMP, so I don't know. But I had expected that you need to specify target on your parallel for pragma. I.e., something like

#pragma omp target teams distribute parallel for

Off course I have #pragma omp target map(to:a[0:n]) map(to:b[0:n]) map(tofrom:c[0:n]) #pragma omp parallel for simd before my loop. Do you have any ideas where it can come from ? Thank's for your help !

I have never played with OpenMP, so I don't know. But I had expected that you need to specify target on your parallel for pragma. I.e., something like

#pragma omp target teams distribute parallel for

So, I look at this and I don't find why it doesn't work on my machine. I used nvc or icx for now and it work well with openmp. If someone who read this have an idea I will be happy to know it. @kristerw Thanks for your work and your help

hello
i installed gcc+nvptx using your script.
gcc version 11.2 was installed normally, and I tried to compile a simple openACC FORTRAN program with gfortran.
but i get following errors:

gfortran -O3 -fopenacc -ffree-line-length-none -foffload=nvptx-none -c global_parameters_mod.F90
gfortran -O3 -fopenacc -ffree-line-length-none -foffload=nvptx-none -c argument_mod.F90
gfortran -O3 -fopenacc -ffree-line-length-none -foffload=nvptx-none -c kind_params_mod.f90
gfortran -O3 -fopenacc -ffree-line-length-none -foffload=nvptx-none -c gocean_mod.F90
gfortran -O3 -fopenacc -ffree-line-length-none -foffload=nvptx-none -c region_mod.f90
gfortran -O3 -fopenacc -ffree-line-length-none -foffload=nvptx-none -c grid_mod.f90
gfortran -O3 -fopenacc -ffree-line-length-none -foffload=nvptx-none -c kernel_mod.F90
gfortran -O3 -fopenacc -ffree-line-length-none -foffload=nvptx-none -c halo_mod.f90
gfortran -O3 -fopenacc -ffree-line-length-none -foffload=nvptx-none -c field_mod.f90
ar rv gocean_api.a argument_mod.o kind_params_mod.o global_parameters_mod.o gocean_mod.o grid_mod.o kernel_mod.o halo_mod.o field_mod.o region_mod.o
ar: creating gocean_api.a
a - argument_mod.o
a - kind_params_mod.o
a - global_parameters_mod.o
a - gocean_mod.o
a - grid_mod.o
a - kernel_mod.o
a - halo_mod.o
a - field_mod.o
a - region_mod.o
make[2]: Leaving directory '/data1/home/synim/test/nemolite2d_edata/api_v1.0'
gfortran -O3 -fopenacc -ffree-line-length-none -foffload=nvptx-none -I../../api_v1.0 -I/src -c nemolite2d.f90
f951: Warning: Nonexistent include directory ‘/src’ [-Wmissing-include-dirs]
gfortran -o nemolite2d.exe nemolite2d.o ../../api_v1.0/gocean_api.a  -lgfortran -lasan -latomic -lgomp
unresolved symbol sin
collect2: error: ld returned 1 exit status
mkoffload: fatal error: x86_64-pc-linux-gnu-accel-nvptx-none-gcc returned 1 exit status
compilation terminated.
lto-wrapper: fatal error: /data1/home/synim/offload/install/libexec/gcc/x86_64-pc-linux-gnu/11.3.1//accel/nvptx-none/mkoffload returned 1 exit status
compilation terminated.
/opt/rh/devtoolset-7/root/usr/bin/ld: error: lto-wrapper failed
collect2: error: ld returned 1 exit status
make[1]: *** [Makefile:40: nemolite2d.exe] Error 1
make[1]: Leaving directory '/data1/home/synim/test/nemolite2d_edata/nemolite2d_serial/nemolite2d_orig'
make: *** [Makefile:24: nemolite2d] Error 2

Anyone who knows about this issue, please help.
thanks.

I have not tried offloading since I wrote the blog post 5 years ago, so my knowledge is a bit rusty... But I think you need to add -foffload=-lm when linking nemolite2d.exe. I also thought you would need -foffload=nvptx-none when linking, but the error message mentions nvptx, so maybe that isn't necessary (but it would not hurt... :) )

Thanks to your advice, the problem has been resolved.
Thank you very much.