[Offload] Offload to NVIDIA GPUs fails with CUDA 13.0 or newer (LLVM 20.1.8)

[Thyre]

I'm still trying to figure out what goes wrong exactly, but maybe someone has an idea.

With LLVM 20.1.8, built with a bootstrapped build with EasyBuild, building any program using OpenMP offload via e.g. -fopenmp --offload-arch=sm_75 fails when running the application with:

omptarget error: Consult https://openmp.llvm.org/design/Runtimes.html for debugging options.
omptarget error: No images found compatible with the installed hardware. [1]    14988 segmentation fault (core dumped)  OMP_TARGET_OFFLOAD=mandatory ./zaxpy

Looking closer with GDB, the stack trace looks like this:

Using host libthread_db library "/usr/lib/x86_64-linux-gnu/libthread_db.so.1".
omptarget error: Consult https://openmp.llvm.org/design/Runtimes.html for debugging options.
omptarget error: No images found compatible with the installed hardware. 
Program received signal SIGSEGV, Segmentation fault.
0x00007ffff5286ccf in llvm::object::ELFObjectFileBase::getNVPTXCPUName() const () from /opt/EasyBuild/apps/software/LLVM/20.1.8-GCCcore-14.3.0/lib/libLLVM.so.20.1
(gdb) bt
#0  0x00007ffff5286ccf in llvm::object::ELFObjectFileBase::getNVPTXCPUName() const () from /opt/EasyBuild/apps/software/LLVM/20.1.8-GCCcore-14.3.0/lib/libLLVM.so.20.1
#1  0x00007ffff5286c53 in llvm::object::ELFObjectFileBase::tryGetCPUName() const () from /opt/EasyBuild/apps/software/LLVM/20.1.8-GCCcore-14.3.0/lib/libLLVM.so.20.1
#2  0x00007ffff7a9cca1 in handleTargetOutcome(bool, ident_t*) () from /opt/EasyBuild/apps/software/LLVM/20.1.8-GCCcore-14.3.0/lib/x86_64-unknown-linux-gnu/libomptarget.so.20.1
#3  0x00007ffff7a97f43 in checkDevice(long&, ident_t*) () from /opt/EasyBuild/apps/software/LLVM/20.1.8-GCCcore-14.3.0/lib/x86_64-unknown-linux-gnu/libomptarget.so.20.1
#4  0x00007ffff7a984e0 in void targetData<AsyncInfoTy>(ident_t*, long, int, void**, void**, long*, long*, void**, void**, int (*)(ident_t*, DeviceTy&, int, void**, void**, long*, long*, void**, void**, AsyncInfoTy&, bool), char const*, char const*) ()
   from /opt/EasyBuild/apps/software/LLVM/20.1.8-GCCcore-14.3.0/lib/x86_64-unknown-linux-gnu/libomptarget.so.20.1
#5  0x00007ffff7a980c4 in __tgt_target_data_begin_mapper () from /opt/EasyBuild/apps/software/LLVM/20.1.8-GCCcore-14.3.0/lib/x86_64-unknown-linux-gnu/libomptarget.so.20.1
#6  0x000055555555ae7f in main ()

Testing CUDA 12.9.1 or earlier, everything looks okay. It seems to only affect CUDA 13.0.0 and 13.0.1 so far.
I haven't tried LLVM 21.1.1 yet, mostly due to the only machine I'm able to test this with taking quite long to build LLVM with.

Its also worth noting that one can build the application with an older CUDA and then run with the newer one. The other way around also fails. Maybe some changes in between these major version causes issues. There were the announced ELF visibility and linkage changes, but as far as I understand, this only affects nvcc. The driver itself should be recent enough (580.65.06).

LLVM 20.1.8 is particularly interesting because of e.g. Numba supporting that particular version soon, while CUDA 13 is interesting for better support of recent GPUs. Mixing LLVM versions would be a noticeable inconvenience.

I'll now try to get a version of LLVM 21 built for cross-checking. Maybe the issue is already resolved and I just haven't found the correct PR for that yet.

[llvmbot]

@llvm/issue-subscribers-offload

Author: Jan André Reuter (Thyre)

I'm still trying to figure out what goes wrong exactly, but maybe someone has an idea.

With LLVM 20.1.8, built with a bootstrapped build with EasyBuild, building any program using OpenMP offload via e.g. -fopenmp --offload-arch=sm_75 fails when running the application with:

omptarget error: Consult https://openmp.llvm.org/design/Runtimes.html for debugging options.
omptarget error: No images found compatible with the installed hardware. [1]    14988 segmentation fault (core dumped)  OMP_TARGET_OFFLOAD=mandatory ./zaxpy

Looking closer with GDB, the stack trace looks like this:

Using host libthread_db library "/usr/lib/x86_64-linux-gnu/libthread_db.so.1".
omptarget error: Consult https://openmp.llvm.org/design/Runtimes.html for debugging options.
omptarget error: No images found compatible with the installed hardware. 
Program received signal SIGSEGV, Segmentation fault.
0x00007ffff5286ccf in llvm::object::ELFObjectFileBase::getNVPTXCPUName() const () from /opt/EasyBuild/apps/software/LLVM/20.1.8-GCCcore-14.3.0/lib/libLLVM.so.20.1
(gdb) bt
#<!-- -->0  0x00007ffff5286ccf in llvm::object::ELFObjectFileBase::getNVPTXCPUName() const () from /opt/EasyBuild/apps/software/LLVM/20.1.8-GCCcore-14.3.0/lib/libLLVM.so.20.1
#<!-- -->1  0x00007ffff5286c53 in llvm::object::ELFObjectFileBase::tryGetCPUName() const () from /opt/EasyBuild/apps/software/LLVM/20.1.8-GCCcore-14.3.0/lib/libLLVM.so.20.1
#<!-- -->2  0x00007ffff7a9cca1 in handleTargetOutcome(bool, ident_t*) () from /opt/EasyBuild/apps/software/LLVM/20.1.8-GCCcore-14.3.0/lib/x86_64-unknown-linux-gnu/libomptarget.so.20.1
#<!-- -->3  0x00007ffff7a97f43 in checkDevice(long&amp;, ident_t*) () from /opt/EasyBuild/apps/software/LLVM/20.1.8-GCCcore-14.3.0/lib/x86_64-unknown-linux-gnu/libomptarget.so.20.1
#<!-- -->4  0x00007ffff7a984e0 in void targetData&lt;AsyncInfoTy&gt;(ident_t*, long, int, void**, void**, long*, long*, void**, void**, int (*)(ident_t*, DeviceTy&amp;, int, void**, void**, long*, long*, void**, void**, AsyncInfoTy&amp;, bool), char const*, char const*) ()
   from /opt/EasyBuild/apps/software/LLVM/20.1.8-GCCcore-14.3.0/lib/x86_64-unknown-linux-gnu/libomptarget.so.20.1
#<!-- -->5  0x00007ffff7a980c4 in __tgt_target_data_begin_mapper () from /opt/EasyBuild/apps/software/LLVM/20.1.8-GCCcore-14.3.0/lib/x86_64-unknown-linux-gnu/libomptarget.so.20.1
#<!-- -->6  0x000055555555ae7f in main ()

Testing CUDA 12.9.1 or earlier, everything looks okay. It seems to only affect CUDA 13.0.0 and 13.0.1 so far.
I haven't tried LLVM 21.1.1 yet, mostly due to the only machine I'm able to test this with taking quite long to build LLVM with.

Its also worth noting that one can build the application with an older CUDA and then run with the newer one. The other way around also fails. Maybe some changes in between these major version causes issues. There were the announced ELF visibility and linkage changes, but as far as I understand, this only affects nvcc. The driver itself should be recent enough (580.65.06).

LLVM 20.1.8 is particularly interesting because of e.g. Numba supporting that particular version soon, while CUDA 13 is interesting for better support of recent GPUs. Mixing LLVM versions would be a noticeable inconvenience.

I'll now try to get a version of LLVM 21 built for cross-checking. Maybe the issue is already resolved and I just haven't found the correct PR for that yet.

[jhuber6]

Interesting. Think you could tell me what llvm-readelf -h says on your image? You should be able to get at it either with -save-temps or by taking it from the .llvm.offloading section via clang-offload-packager a.out --image=file=out.cubin

[jhuber6]

Also do you know if it works upstream, I made some changes there recently.

[Thyre]

Interesting. Think you could tell me what llvm-readelf -h says on your image? You should be able to get at it either with -save-temps or by taking it from the .llvm.offloading section via clang-offload-packager a.out --image=file=out.cubin

Here's the output for the two .cubin files with llvm-readelf -h and the .img file:

$ llvm-readelf -h zaxpy.nvptx64.sm_75.cubin 
ELF Header:
  Magic:   7f 45 4c 46 02 01 01 41 08 00 00 00 00 00 00 00
  Class:                             ELF64
  Data:                              2's complement, little endian
  Version:                           1 (current)
  OS/ABI:                            41
  ABI Version:                       8
  Type:                              REL (Relocatable file)
  Machine:                           NVIDIA CUDA architecture
  Version:                           0x1
  Entry point address:               0x0
  Start of program headers:          0 (bytes into file)
  Start of section headers:          912 (bytes into file)
  Flags:                             0x6004B04
  Size of this header:               64 (bytes)
  Size of program headers:           56 (bytes)
  Number of program headers:         0
  Size of section headers:           64 (bytes)
  Number of section headers:         10
  Section header string table index: 1
$ llvm-readelf -h zaxpy-host-x86_64-unknown-linux-gnu-nvptx64-nvidia-cuda-sm_75.cubin 
ELF Header:
  Magic:   7f 45 4c 46 02 01 01 41 08 00 00 00 00 00 00 00
  Class:                             ELF64
  Data:                              2's complement, little endian
  Version:                           1 (current)
  OS/ABI:                            41
  ABI Version:                       8
  Type:                              REL (Relocatable file)
  Machine:                           NVIDIA CUDA architecture
  Version:                           0x1
  Entry point address:               0x0
  Start of program headers:          0 (bytes into file)
  Start of section headers:          29536 (bytes into file)
  Flags:                             0x6004B04
  Size of this header:               64 (bytes)
  Size of program headers:           56 (bytes)
  Number of program headers:         0
  Size of section headers:           64 (bytes)
  Number of section headers:         38
  Section header string table index: 1
$ llvm-readelf -h zaxpy.nvptx64.sm_75.img 
ELF Header:
  Magic:   7f 45 4c 46 02 01 01 41 08 00 00 00 00 00 00 00
  Class:                             ELF64
  Data:                              2's complement, little endian
  Version:                           1 (current)
  OS/ABI:                            41
  ABI Version:                       8
  Type:                              EXEC (Executable file)
  Machine:                           NVIDIA CUDA architecture
  Version:                           0x1
  Entry point address:               0x0
  Start of program headers:          19040 (bytes into file)
  Start of section headers:          17504 (bytes into file)
  Flags:                             0x6004B04
  Size of this header:               64 (bytes)
  Size of program headers:           56 (bytes)
  Number of program headers:         4
  Size of section headers:           64 (bytes)
  Number of section headers:         24
  Section header string table index: 1

Also do you know if it works upstream, I made some changes there recently.

I started a build, but that failed with OOM. So I need to start another build this evening or tomorrow morning on one of our nodes where the CUDA driver is new enough. Sorry 😕

[jhuber6]

Looks like an old LLVM readelf, maybe we need to backport #149534?

[Thyre]

I can try to check if the changes of that PR can be applied to LLVM 20.1.8. Would be great if that's sufficient.
I'll put that on my list for tomorrow morning, as it's getting late here.

[jhuber6]

Alright, unfortunately I can't easily test CUDA 13.0 because it's not in my system's packages yet. Can you give me the ELF for this simple program? I might be able to hack in the image so I can test it.

int main() {
  #pragma omp target
  ;
}

clang main.c --offload-arch=sm_89 -fopenmp -O2

[Thyre]

Here's a zip archive with the resulting binary & the temp files (if needed): reproducer.zip

I'll now try to build LLVM 21.1.1 to cross-check (and trunk if LLVM 21 doesn't resolve this yet).

[Thyre]

LLVM 21.1.1 fails the same way (which makes sense as b53be5f is not included there, if I haven't missed something).

Will now try trunk

[jhuber6]

Let me know if it works with that PR applied, as far as I can tell the CUDA 13.0 update applied the new ABI changes to all the SMs wheras previously it was only used for the ones greater than 100.