<div dir="ltr"><div>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.</div><div>Ye<br></div><div><div><div dir="ltr" class="gmail_signature" data-smartmail="gmail_signature"><div dir="ltr"><div><div dir="ltr">===================<br>
Ye Luo, Ph.D.<br>Computational Science Division & Leadership Computing Facility<br>
Argonne National Laboratory</div></div></div></div></div><br></div></div><br><div class="gmail_quote"><div dir="ltr" class="gmail_attr">On Tue, Jan 4, 2022 at 5:59 PM pramod kumbhar via Openmp-dev <<a href="mailto:openmp-dev@lists.llvm.org">openmp-dev@lists.llvm.org</a>> wrote:<br></div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex"><div dir="ltr">Dear All,<div><br></div><div>I have two questions regarding the usage of global variables with OpenMP offload:</div><div><br></div><div>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:</div><div><br></div><blockquote style="margin:0px 0px 0px 40px;border:medium none;padding:0px"><div><font color="#20124d">$ cat test.sh</font></div><div><font color="#20124d"><br></font></div><div><font color="#20124d">CXX=clang++</font></div><div><font color="#20124d">CXXFLAGS="-fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -g -O2"</font></div><div><font color="#20124d">${CXX} ${CXXFLAGS} -c test.cpp</font></div><div><font color="#20124d">ar cq libtest.a test.o</font></div><div><font color="#20124d">${CXX} ${CXXFLAGS} -o test1 main.cpp -L. -ltest</font></div><div><font color="#20124d">${CXX} ${CXXFLAGS} -o test2 main.cpp test.o</font></div><div><font color="#20124d"><br></font></div><div><font color="#0b5394">$ cat test.cpp</font></div><div><font color="#0b5394"><br>#pragma omp declare target<br>int y;<br>#pragma omp end declare target<br><br>int test() {<br> y = 24;<br> #pragma omp target update to(y)<br> y = 42;<br><br> int x;<br> #pragma omp target map(from:x)<br> {<br> x = y;<br> }<br> return x;<br>}</font></div><div><font color="#0b5394"><br>$ cat main.cpp<br>extern int test();<br><br>int main() {<br> return test();<br>}</font><font color="#20124d"><br></font></div></blockquote><div><br></div>Running the ./<i>test2</i> works as expected as I am not using static library but the ./test1 fails with an error shown below:<div><br><blockquote style="margin:0px 0px 0px 40px;border:medium none;padding:0px"><div><font color="#351c75">$ ./test2</font></div><div><font color="#351c75">$ echo $?</font></div><div><font color="#351c75">24</font></div><div><font color="#351c75">$ ./test1</font></div><div><font color="#351c75">CUDA error: Loading global 'y' Failed</font></div><div><font color="#351c75">CUDA error: named symbol not found</font></div><div><font color="#351c75">Libomptarget error: Unable to generate entries table for device id 0.</font></div><div><font color="#351c75">Libomptarget error: Failed to init globals on device 0</font></div><div><font color="#351c75">Libomptarget error: Run with LIBOMPTARGET_INFO=4 to dump host-target pointer mappings.</font></div><div><font color="#351c75">test.cpp:7:3: Libomptarget fatal error 1: failure of target construct while offloading is mandatory</font></div><div><font color="#351c75">Aborted</font></div></blockquote><div><blockquote style="margin:0px 0px 0px 40px;border:medium none;padding:0px"><div><br></div></blockquote>Is this expected behaviour? Is there any workaround? I have tested this with PGI/NVHPC compiler and it works there.</div><div><br></div><div>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:</div><div><br></div><blockquote style="margin:0px 0px 0px 40px;border:medium none;padding:0px"><div><font color="#20124d">$ cat main.cpp</font></div><div><font color="#20124d">extern int test();</font></div><div><font color="#20124d"><br></font></div><div><font color="#20124d">#include <cstdio></font></div><div><font color="#20124d"><br></font></div><div><font color="#20124d">#pragma omp declare target</font></div><div><font color="#20124d">extern int y;</font></div><div><font color="#20124d">#pragma omp end declare target</font></div><div><font color="#20124d"><br></font></div><div><font color="#20124d">int main() {</font></div><div><font color="#20124d"> #pragma omp target teams distribute parallel for</font></div><div><font color="#20124d"> for(int i=0; i<5; i++) {</font></div><div><font color="#20124d"> printf("--> %d \n", y + i);</font></div><div><font color="#20124d"> }</font></div><div><font color="#20124d"> return test();</font></div><div><font color="#20124d">}</font></div><div><font color="#20124d"><br></font></div></blockquote><font color="#20124d">This now fails to compile with a library:</font><div><font color="#20124d"><br></font></div><blockquote style="margin:0px 0px 0px 40px;border:medium none;padding:0px"><div><font color="#351c75">+ clang++ -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -g -O2 -c test.cpp</font></div><div><font color="#351c75">+ ar cq libtest.a test.o</font></div><div><font color="#351c75">+ clang++ -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -g -O2 -o test1 main.cpp -L. -ltest</font></div><div><font color="#351c75">nvlink error : Undefined reference to 'y' in '/gpfs/<a href="http://bbp.cscs.ch/ssd/slurmTmpFS/kumbhar/151151/main-9ea2c6.cubin" target="_blank">bbp.cscs.ch/ssd/slurmTmpFS/kumbhar/151151/main-9ea2c6.cubin</a>'</font></div><div><font color="#351c75">clang-13: error: nvlink command failed with exit code 255 (use -v to see invocation)</font></div><div><font color="#351c75">clang version 13.0.0</font></div><div><font color="#351c75">Target: x86_64-unknown-linux-gnu</font></div><div><font color="#351c75">Thread model: posix</font></div><div><font color="#351c75">InstalledDir: /gpfs/<a href="http://bbp.cscs.ch/ssd/apps/bsd/2021-11/stage_externals/install_gcc-11.2.0-skylake/llvm-13.0.0-lvcrm6/bin" target="_blank">bbp.cscs.ch/ssd/apps/bsd/2021-11/stage_externals/install_gcc-11.2.0-skylake/llvm-13.0.0-lvcrm6/bin</a></font></div><div><font color="#351c75"><br></font></div></blockquote><font color="#000000"><span>In a big application, we build & ship libraries and it's not easy/convenient to use objects for linking. </span></font><span style="color:rgb(0,0,0)">Do you have any recommendations to solve/workaround issues for this use case?</span><div><div><br></div><div>Thank you very much!</div><div><br></div><div>Regards,</div><div>Pramod Kumbhar</div></div></div></div>
_______________________________________________<br>
Openmp-dev mailing list<br>
<a href="mailto:Openmp-dev@lists.llvm.org" target="_blank">Openmp-dev@lists.llvm.org</a><br>
<a href="https://lists.llvm.org/cgi-bin/mailman/listinfo/openmp-dev" rel="noreferrer" target="_blank">https://lists.llvm.org/cgi-bin/mailman/listinfo/openmp-dev</a><br>
</blockquote></div>