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

[Ye Luo via Openmp-dev]

As of today, linking static libraries with device code in it remains
problematic. If your project uses CMake, use object target to link object
files directly.
Ye
===================
Ye Luo, Ph.D.
Computational Science Division & Leadership Computing Facility
Argonne National Laboratory


On Tue, Jan 4, 2022 at 5:59 PM pramod kumbhar via Openmp-dev <
openmp-dev at lists.llvm.org> wrote:

> 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
> _______________________________________________
> Openmp-dev mailing list
> Openmp-dev at lists.llvm.org
> https://lists.llvm.org/cgi-bin/mailman/listinfo/openmp-dev
>
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/openmp-dev/attachments/20220104/9833291b/attachment.html>