title: nvlink, relocatable device code and static device libraries
body: <p>While investigating some issues with relocatable device code, I stumbled upon something I don't quite understand.</p>
  
  <p>This is a use case for what is pictured on <a href="http://on-demand.gputechconf.com/gtc-express/2012/presentations/gpu-object-linking.pdf#6" rel="nofollow noreferrer">slide 6</a>. I used an <a href="https://stackoverflow.com/a/17287863/1043187">answer of Robert Crovella</a> as a basis for a repro code. The idea is that we have some relocatable device code compiled into a static library (e.g. some math/toolbox library), and we want to use some functions of that precompiled library into another device library of our program:</p>
  
  
  
  <pre><code>libutil.a ---&gt; libtest.so ---&gt; test_pgm
  </code></pre>
  
  <p>Let's say that this external library contains the following function:</p>
  
  
  
  <pre><code>__device__ int my_square (int a);
  </code></pre>
  
  <p><code>libutil.a</code> was for instance generated as follow (in another project):</p>
  
  <pre class="lang-sh prettyprint-override"><code>nvcc ${NVCC_FLAGS} -dc util.cu
  nvcc ${NVCC_FLAGS} -dlink util.o -o util_dlink.o
  nvcc ${NVCC_FLAGS} -lib util_dlink.o util.o -o libutil.a
  </code></pre>
  
  <p>Then, in our project, to generate <code>libtest.so</code>:</p>
  
  <pre class="lang-sh prettyprint-override"><code>nvcc ${NVCC_FLAGS} -dc test.cu
  nvcc ${NVCC_FLAGS} -dlink test.o libutil.a -o test_dlink.o
  g++ -shared -Wl,-soname,libtest.so -o libtest.so test.o test_dlink.o libutil.a -L${CUDA_LIBDIR} -lcudart
  </code></pre>
  
  <p>But I get the following error when generating <code>test_dlink.o</code>:</p>
  
  <pre><code>nvlink error   : Undefined reference to '_Z9my_squarei' in 'test.o'
  </code></pre>
  
  <p>The linker does not find our dummy <code>my_square(int)</code> function. If we instead use (assuming we had access to <code>util.o</code>):</p>
  
  <pre><code>nvcc ${NVCC_FLAGS} -dlink test.o util.o -o test_dlink.o
  </code></pre>
  
  <p>The linker succeeds and everything works fine after.</p>
  
  <p>Investigating further:</p>
  
  <pre><code>$ nm -C libutil.a
  
  util_dlink.o:
                   U atexit
                   U __cudaRegisterFatBinary
  0000000000000015 T __cudaRegisterLinkedBinary_39_tmpxft_0000106a_00000000_6_util_cpp1_ii_843d693d
    ...
  
  util.o:
                   U __cudaInitModule
                   U __cudaRegisterLinkedBinary_39_tmpxft_0000106a_00000000_6_util_cpp1_ii_843d693d
    ...
  0000000000000015 T my_square(int)
    ...
  </code></pre>
  
  <p>The symbol is there in the archived <code>util.o</code>, but <code>nvlink</code> (called by <code>nvcc</code>) does not seem to find it. Why is that? According to the <a href="http://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#libraries" rel="nofollow noreferrer">official documentation</a>:</p>
  
  <blockquote>
    <p>The device linker has the ability to read the static host library
    formats (.a on Linux and Mac, .lib on Windows).</p>
  </blockquote>
  
  <p>We could of course extract the object file and link with it:</p>
  
  <pre class="lang-sh prettyprint-override"><code>ar x libutil.a `ar t libutil.a | grep -v "dlink"`
  nvcc ${NVCC_FLAGS} -dlink test.o util.o -o test_dlink.o
  </code></pre>
  
  <p>But this does not feel like the expected solution... So what am I missing here? Another <code>nvcc</code> option that solves that? Is there an error when generating <code>libutil.a</code> and/or <code>libtest.so</code>?</p>
  
  <p>Note that this was tested with CUDA 6.5 on Arch Linux.</p>
  
  <h2>EDIT: fixed repro code with commented lines</h2>
  
  <h3>Makefile</h3>
  
  <pre><code>NVCC_FLAGS=-m64 -arch=sm_20 -Xcompiler '-fPIC'
  CUDA_LIBDIR=${CUDA_HOME}/lib64
  
  testmain : main.cpp libtest.so
      g++ -c main.cpp
      g++ -o testmain -L. -ldl -Wl,-rpath,. -ltest -L${CUDA_LIBDIR} -lcudart main.o
  
  libutil.a : util.cu util.cuh
      nvcc ${NVCC_FLAGS} -dc util.cu
      # ---&gt; FOLLOWING LINES THAT WERE WRONG &lt;---
      # nvcc ${NVCC_FLAGS} -dlink util.o -o util_dlink.o
      # nvcc ${NVCC_FLAGS} -lib util.o util_dlink.o -o libutil.a
      # INSTEAD:
      nvcc ${NVCC_FLAGS} -lib util.o -o libutil.a
      # Assuming util is an external library, so util.o is not available
      rm util.o
  
  libtest.so : test.cu test.h libutil.a util.cuh
      nvcc ${NVCC_FLAGS} -dc test.cu
      # Use NVCC for device linking + G++
      nvcc -v ${NVCC_FLAGS} -dlink test.o libutil.a -o test_dlink.o
      g++ -shared -o libtest.so test.o test_dlink.o libutil.a -L${CUDA_LIBDIR} -lcudart
      # Or let NVCC generate the shared library
      #nvcc -v ${NVCC_FLAGS} -shared -L. -lutil test.o -o libtest.so
  
  clean :
      rm -f testmain *.o *.a *.so
  </code></pre>
  
  <h3>test.h</h3>
  
  <pre><code>#ifndef TEST_H
  # define TEST_H
  
  int my_test_func();
  
  #endif //! TEST_H
  </code></pre>
  
  <h3>test.cu</h3>
  
  <pre><code>#include &lt;stdio.h&gt;
  
  #include "test.h"
  #include "util.cuh"
  
  #define DSIZE 1024
  #define DVAL 10
  #define SQVAL 3
  #define nTPB 256
  
  #define cudaCheckErrors(msg)                             \
    do {                                                   \
      cudaError_t __err = cudaGetLastError();              \
      if (__err != cudaSuccess) {                          \
        fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err),            \
                __FILE__, __LINE__);                       \
        fprintf(stderr, "*** FAILED - ABORTING\n");        \
        exit(1);                                           \
      }                                                    \
    } while (0)
  
  __global__ void my_kernel(int *data){
    int idx = threadIdx.x + (blockDim.x *blockIdx.x);
    if (idx &lt; DSIZE) data[idx] =+ DVAL + my_square (SQVAL);
  }
  
  int my_test_func()
  {
    int *d_data, *h_data;
    h_data = (int *) malloc(DSIZE * sizeof(int));
    if (h_data == 0) {printf("malloc fail\n"); exit(1);}
    cudaMalloc((void **)&amp;d_data, DSIZE * sizeof(int));
    cudaCheckErrors("cudaMalloc fail");
    for (int i = 0; i &lt; DSIZE; i++) h_data[i] = 0;
    cudaMemcpy(d_data, h_data, DSIZE * sizeof(int), cudaMemcpyHostToDevice);
    cudaCheckErrors("cudaMemcpy fail");
    my_kernel&lt;&lt;&lt;((DSIZE+nTPB-1)/nTPB), nTPB&gt;&gt;&gt;(d_data);
    cudaDeviceSynchronize();
    cudaCheckErrors("kernel");
    cudaMemcpy(h_data, d_data, DSIZE * sizeof(int), cudaMemcpyDeviceToHost);
    cudaCheckErrors("cudaMemcpy 2");
    for (int i = 0; i &lt; DSIZE; i++)
      if (h_data[i] != DVAL + SQVAL*SQVAL)
      {
        printf("Results check failed at offset %d, data was: %d, should be %d\n",
               i, h_data[i], DVAL);
        exit(1);
      }
    printf("Results check passed!\n");
    return 0;
  }
  </code></pre>
  
  <h3>util.cuh</h3>
  
  <pre><code>#ifndef UTIL_CUH
  # define UTIL_CUH
  
  __device__ int my_square (int a);
  
  #endif //! UTIL_CUH
  </code></pre>
  
  <h3>util.cu</h3>
  
  <pre><code>#include "util.cuh"
  
  __device__ int my_square (int a)
  {
    return a * a;
  }
  </code></pre>
  
  <h3>main.cpp</h3>
  
  <pre><code>#include "test.h"
  
  int main()
  {
    my_test_func();
    return 0;
  }
  </code></pre>