[OpenMP] Offloading returns incorrect value when compiled from a different file

[jhuber6]

The following offloading program returns an incorrect value when run on my V100.

/// impl.c
#pragma omp begin declare target
int x = 1;

int function() {
#pragma omp parallel for
  for (int i = 0; i < 128; ++i) {
    x = 0;
  }
  return 0;
}
#pragma omp end declare target

// kernel.c
int function();

int main() {
  int x = 0;
#pragma omp target map(from : x)
  x = function();
  
  return x;
}

When compiled together we get incorrect results using regular compilation.

3n4:iguazu ~/llvm/llvm-project/build> clang impl.c kernel.c -fopenmp -fopenmp-targets=nvptx64 -O3
3n4:iguazu ~/llvm/llvm-project/build> ./a.out; echo $?
33

However, using LTO or manually linking the static library this seems fine.

3n4:iguazu ~/llvm/llvm-project/build> clang impl.c kernel.c -fopenmp -fopenmp-targets=nvptx64 -O3 -foffload-lto
3n4:iguazu ~/llvm/llvm-project/build> ./a.out; echo $?
0

This is most likely caused by a variable that is not shred correctly in the bitcode version of the runtime. We can try to identify this variable or move forward making LTO default / always using the static library.

[llvmbot]

@llvm/issue-subscribers-openmp