[OpenMP] Compilation error at O0 for OpenMP GPU offloading with NVIDIA GPU and LLVM 20

[ouankou]

Given the following OpenMP C code, at O0 there is a compilation error. O1-3 works fine. The issue occurs in LLVM 20 but not the main branch (LLVM 21).

#include <stdio.h>

void foo() {
  int sum = 10000;
#pragma omp target teams distribute parallel for map(tofrom: sum) num_teams(256) num_threads(1024) reduction(+ : sum)
  for (int i = 1; i <= 100; i++)
    sum += i;

  printf("sum = %d\n", sum);
}

int main(int argc, char **argv) {

  foo();

  return 0;
}

$ clang -v
clang version 20.1.2 (https://github.com/llvm/llvm-project 3f957cc67cff4e337f56fa2dbfdb037d3a997baf)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /home/ouankou/Projects/llvm-20/llvm_install/bin
Found candidate GCC installation: /usr/lib/gcc/x86_64-linux-gnu/11
Found candidate GCC installation: /usr/lib/gcc/x86_64-linux-gnu/12
Selected GCC installation: /usr/lib/gcc/x86_64-linux-gnu/12
Candidate multilib: .;@m64
Selected multilib: .;@m64
Found CUDA installation: /opt/nvidia/hpc_sdk/Linux_x86_64/24.3/cuda/12.3, version 12.3

$ clang -fopenmp -fopenmp-targets=nvptx64 -O0 test.c
nvlink error   : Undefined reference to 'IsSPMDMode' in '/tmp/a.out.nvptx64.sm_86-a57cc1-8872c3.cubin'
nvlink error   : Undefined reference to '_ZN4ompx5state9TeamStateE' in '/tmp/a.out.nvptx64.sm_86-a57cc1-8872c3.cubin'
nvlink error   : Undefined reference to '_ZN4ompx5state12ThreadStatesE' in '/tmp/a.out.nvptx64.sm_86-a57cc1-8872c3.cubin'
/home/ouankou/Projects/llvm-20/llvm_install/bin/clang-nvlink-wrapper: error: 'nvlink' failed
clang: error: nvlink command failed with exit code 1 (use -v to see invocation)
/home/ouankou/Projects/llvm-20/llvm_install/bin/clang-linker-wrapper: error: 'clang' failed
clang: error: linker command failed with exit code 1 (use -v to see invocation)

The setup is LLVM 20 built from source, NVIDIA HPC SDK 24.3 (CUDA 12.3), NVIDIA RTX A3000 Laptop GPU (sm86). The full error log is attached. Please let me know if any other information is needed. Thanks!

error.log

[llvmbot]

@llvm/issue-subscribers-openmp

Author: Anjia Wang (ouankou)

Given the following OpenMP C code, at O0 there is a compilation error. O1-3 works fine. The issue occurs in LLVM 20 but not the main branch (LLVM 21). ```c #include <stdio.h>

void foo() {
int sum = 10000;
#pragma omp target teams distribute parallel for map(tofrom: sum) num_teams(256) num_threads(1024) reduction(+ : sum)
for (int i = 1; i <= 100; i++)
sum += i;

printf("sum = %d\n", sum);
}

int main(int argc, char **argv) {

foo();

return 0;
}

```bash
$ clang -v
clang version 20.1.2 (https://github.com/llvm/llvm-project 3f957cc67cff4e337f56fa2dbfdb037d3a997baf)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /home/ouankou/Projects/llvm-20/llvm_install/bin
Found candidate GCC installation: /usr/lib/gcc/x86_64-linux-gnu/11
Found candidate GCC installation: /usr/lib/gcc/x86_64-linux-gnu/12
Selected GCC installation: /usr/lib/gcc/x86_64-linux-gnu/12
Candidate multilib: .;@<!-- -->m64
Selected multilib: .;@<!-- -->m64
Found CUDA installation: /opt/nvidia/hpc_sdk/Linux_x86_64/24.3/cuda/12.3, version 12.3

$ clang -fopenmp -fopenmp-targets=nvptx64 -O0 test.c
nvlink error   : Undefined reference to 'IsSPMDMode' in '/tmp/a.out.nvptx64.sm_86-a57cc1-8872c3.cubin'
nvlink error   : Undefined reference to '_ZN4ompx5state9TeamStateE' in '/tmp/a.out.nvptx64.sm_86-a57cc1-8872c3.cubin'
nvlink error   : Undefined reference to '_ZN4ompx5state12ThreadStatesE' in '/tmp/a.out.nvptx64.sm_86-a57cc1-8872c3.cubin'
/home/ouankou/Projects/llvm-20/llvm_install/bin/clang-nvlink-wrapper: error: 'nvlink' failed
clang: error: nvlink command failed with exit code 1 (use -v to see invocation)
/home/ouankou/Projects/llvm-20/llvm_install/bin/clang-linker-wrapper: error: 'clang' failed
clang: error: linker command failed with exit code 1 (use -v to see invocation)

The setup is LLVM 20 built from source, NVIDIA HPC SDK 24.3 (CUDA 12.3), NVIDIA RTX A3000 Laptop GPU (sm86). The full error log is attached. Please let me know if any other information is needed. Thanks!

error.log

[jhuber6]

I don't have any good guesses for why this works on the main branch but not in the 20 release. The errors suggest it's some missing symbols, those specific ones always problematic due to the weird linking situation we have. Does it work if you pass -foffload-lto? There's still some time to backport things, but honestly if it's working upstream I'm not sure how high of a priority it is to backport unless you bisect it to the specific commit that fixed it.

[ouankou]

Thanks for the quick response! -foffload-lto works.
clang -fopenmp -fopenmp-targets=nvptx64 -foffload-lto -O0 test.c built and ran the program without any issues.

[jhuber6]

What I figured, those symbols are majorly hacked in. You could likely also make it work by passing -nogpulib and then manually linking libomptarget.devicertl.a (Will rework this in 21 so don't get used to the name). Unfortunately I'm not sure if fixing this is a high priority, but it's related to the use of -mlink-builtin-bitcode which is only here for performance reasons. if we get functional thin-LTO for GPUs I will flip the switch and make that the default, but for now this is kind of broken.

[ouankou]

Sure, it's just for personal experiments, so I can switch to the main branch without problems. Appreciate your detailed explanation!