[Openmp-dev] OpenMP offload: Using global variable with a library result in "CUDA error: Loading global 'xxxx' Failed" or "nvlink error : Undefined reference to ..."

[pramod kumbhar via Openmp-dev]

Dear All,

I have two questions regarding the usage of global variables with OpenMP
offload:

1. When I have global variable usage in the code from which I create a
library then I get "CUDA error: Loading global 'x' Failed" error. Here is a
simple reproducer showing the issue:

$ cat test.sh

CXX=clang++
CXXFLAGS="-fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -g -O2"
${CXX} ${CXXFLAGS} -c test.cpp
ar cq libtest.a test.o
${CXX} ${CXXFLAGS} -o test1 main.cpp -L. -ltest
${CXX} ${CXXFLAGS} -o test2 main.cpp test.o

$ cat test.cpp

#pragma omp declare target
int y;
#pragma omp end declare target

int test() {
  y = 24;
  #pragma omp target update to(y)
  y = 42;

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

$ cat main.cpp
extern int test();

int main() {
  return test();
}


Running the ./*test2* works as expected as I am not using static library
but the ./test1 fails with an error shown below:

$ ./test2
$ echo $?
24
$ ./test1
CUDA error: Loading global 'y' Failed
CUDA error: named symbol not found
Libomptarget error: Unable to generate entries table for device id 0.
Libomptarget error: Failed to init globals on device 0
Libomptarget error: Run with LIBOMPTARGET_INFO=4 to dump host-target
pointer mappings.
test.cpp:7:3: Libomptarget fatal error 1: failure of target construct while
offloading is mandatory
Aborted


Is this expected behaviour? Is there any workaround? I have tested this
with PGI/NVHPC compiler and it works there.

2. The second scenario is similar but now I am trying to use a global
variable from the library into the offload region in main.cpp i.e. modified
main.cpp looks as:

$ cat main.cpp
extern int test();

#include <cstdio>

#pragma omp declare target
extern int y;
#pragma omp end declare target

int main() {
  #pragma omp target teams distribute parallel for
  for(int i=0; i<5; i++) {
    printf("--> %d \n", y + i);
  }
  return test();
}

This now fails to compile with a library:

+ clang++ -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -g -O2 -c test.cpp
+ ar cq libtest.a test.o
+ clang++ -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -g -O2 -o test1
main.cpp -L. -ltest
nvlink error   : Undefined reference to 'y' in '/gpfs/
bbp.cscs.ch/ssd/slurmTmpFS/kumbhar/151151/main-9ea2c6.cubin'
clang-13: error: nvlink command failed with exit code 255 (use -v to see
invocation)
clang version 13.0.0
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /gpfs/
bbp.cscs.ch/ssd/apps/bsd/2021-11/stage_externals/install_gcc-11.2.0-skylake/llvm-13.0.0-lvcrm6/bin

In a big application, we build & ship libraries and it's not
easy/convenient to use objects for linking. Do you have any recommendations
to solve/workaround issues for this use case?

Thank you very much!

Regards,
Pramod Kumbhar
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/openmp-dev/attachments/20220105/648b46ca/attachment.html>