<div dir="ltr"><div dir="ltr">Thank you Ye and Joseph for clarifications! I am glad that the issues is already being addressed!<div><br></div><div>Just would like to add one clarification - the issue/question #1 is addressed by switching to shared library but the issue/question #2 (i.e. extern global variables in offload region) remains with the shared library :</div><div><br></div><blockquote style="margin:0 0 0 40px;border:none;padding:0px"><div><font color="#351c75">clang++ -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -g -O2 -fpic -shared test.cpp -o libtest.so</font></div><div><font color="#351c75">clang++ -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -g -O2 -o test3 main.cpp -L. -ltest -Wl,-rpath .</font></div><div><font color="#351c75">nvlink error : Undefined reference to 'y' in '/gpfs/<a href="http://bbp.cscs.ch/ssd/slurmTmpFS/kumbhar/151397/main-1999e6.cubin">bbp.cscs.ch/ssd/slurmTmpFS/kumbhar/151397/main-1999e6.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></blockquote></div><div><br></div><div>I haven't built the PR <a href="https://reviews.llvm.org/D116541" id="gmail-m_-5332740749121448046LPNoLPOWALinkPreview" target="_blank" style="font-family:Calibri,Arial,Helvetica,sans-serif">D116541</a> locally yet but I assume the underlying issue is the same and will be addressed by the same PR?</div><div><br></div><div>Thank you again!</div><br><div class="gmail_quote"><div dir="ltr" class="gmail_attr">On Wed, Jan 5, 2022 at 3:06 AM Huber, Joseph <<a href="mailto:huberjn@ornl.gov">huberjn@ornl.gov</a>> wrote:<br></div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left-width:1px;border-left-style:solid;border-left-color:rgb(204,204,204);padding-left:1ex">
<div dir="ltr">
<div style="font-family:Calibri,Arial,Helvetica,sans-serif;font-size:12pt;color:rgb(0,0,0)">
Kumbhar,<br>
<br>
As ye said, the current support for linking static libraries is a little flaky. I'm currently working on a series of patches to hopefully address this problem that is under review at
<a href="https://reviews.llvm.org/D116541" id="gmail-m_-5332740749121448046LPNoLPOWALinkPreview" target="_blank">https://reviews.llvm.org/D116541</a> and its child revisions. I can currently link and run your tests using my development branch after adding some additional support to search the library
paths for static libraries. This should solve your problem once these patches are ready to land upstream.<br>
<br>
Thanks,<br>
Joseph Huber</div>
<div></div>
<div id="gmail-m_-5332740749121448046appendonsend"></div>
<hr style="display:inline-block;width:98%">
<div id="gmail-m_-5332740749121448046divRplyFwdMsg" dir="ltr"><font face="Calibri, sans-serif" style="font-size:11pt" color="#000000"><b>From:</b> Openmp-dev <<a href="mailto:openmp-dev-bounces@lists.llvm.org" target="_blank">openmp-dev-bounces@lists.llvm.org</a>> on behalf of pramod kumbhar via Openmp-dev <<a href="mailto:openmp-dev@lists.llvm.org" target="_blank">openmp-dev@lists.llvm.org</a>><br>
<b>Sent:</b> Tuesday, January 4, 2022 6:59 PM<br>
<b>To:</b> <a href="mailto:openmp-dev@lists.llvm.org" target="_blank">openmp-dev@lists.llvm.org</a> <<a href="mailto:openmp-dev@lists.llvm.org" target="_blank">openmp-dev@lists.llvm.org</a>><br>
<b>Subject:</b> [EXTERNAL] [Openmp-dev] OpenMP offload: Using global variable with a library result in "CUDA error: Loading global 'xxxx' Failed" or "nvlink error : Undefined reference to ..."</font>
<div> </div>
</div>
<div>
<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: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: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: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: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: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>
</div>
</div>
</blockquote></div></div>