title: 'inline' for __global__ functions to avoid multiple definition error
body: <p>I have a CUDA template library, in which one function is actually <em>not</em> a template, but <em>is</em> defined within a <code>.cuh</code> header. (<code>vector_add_kernel</code> in <code>kernel.cuh</code> below.)</p>
  
  <p>If multiple <code>.cu</code> files include <code>kernel.cuh</code> and call <code>vector_add[_kernel]</code>, it will result in multiple definition errors at link-time. In C++, one can use the <code>inline</code> qualifier to avoid such errors.</p>
  
  <p>However, <code>inline __global__ ...</code> - while <strong>preventing the multiple definition errors</strong> on my system - results in a warning that the <code>inline</code> qualifier has been ignored.</p>
  
  <p><strong>Q:</strong> Is there a better way to avoid the multiple definition error, or a way to suppress this warning only for this function? And is <code>inline __global__</code> even safe, or might other host compilers <em>truly</em> ignore it?</p>
  
  <p>I could simply move the <code>vector_add_kernel</code> to a separate <code>.cu</code> file, but it would be the <em>only</em> non-header file. I could also template <code>vector_add_kernel</code>, but in my library that makes little sense.</p>
  
  <p>A (not-so-minimal, sorry) working example (tested with CUDA 7.0, gcc 4.7.2 on Debian) is below.</p>
  
  <p>To clarify, <code>main.cu</code> is some user's code; <code>lib.cu</code> is some external library <em>not belonging to me</em>; and <code>kernel.cuh</code> is part of my template library. So, both the external <code>lib</code> and the user's <code>main</code> are using my template library, <code>kernel.cuh</code> - but separately.</p>
  
  <p><code>main.cu</code>:</p>
  
  <pre><code>#include "lib.hpp"
  #include "kernel.cuh"
  
  #include &lt;thrust/device_vector.h&gt;
  #include &lt;thrust/host_vector.h&gt;
  
  #include &lt;cstddef&gt;
  #include &lt;cstdlib&gt;
  #include &lt;iostream&gt;
  
  int main(void)
  {
      const size_t N = 1u &lt;&lt; 7;
  
      float* a = (float*) malloc(N * sizeof(float));
      float* b = (float*) malloc(N * sizeof(float));
      float* c = (float*) malloc(N * sizeof(float));
  
      for (int i = 0; i &lt; N; ++i) {
          a[i] = b[i] = 2.0f * i;
      }
  
      lib_vector_add(a, b, c, N);
      for (int i = 0; i &lt; N; ++i) {
          if (c[i] != 2.0f * i + 2.0f * i)
              std::cout &lt;&lt; "Error, lib, element " &lt;&lt; i &lt;&lt; std::endl;
      }
  
      thrust::device_vector&lt;float&gt; d_a(a, a + N);
      thrust::device_vector&lt;float&gt; d_b(b, b + N);
      thrust::device_vector&lt;float&gt; d_c(N);
  
      vector_add(d_a, d_b, d_c);
      thrust::host_vector&lt;float&gt; h_c = d_c;
      for (int i = 0; i &lt; N; ++i) {
          if (h_c[i] != 2.0f * i + 2.0f * i)
              std::cout &lt;&lt; "Error, element " &lt;&lt; i &lt;&lt; std::endl;
      }
  }
  </code></pre>
  
  <p><code>lib.cu</code>,</p>
  
  <pre><code>#include &lt;kernel.cuh&gt;
  
  #include &lt;thrust/host_vector.h&gt;
  #include &lt;thrust/device_vector.h&gt;
  
  void lib_vector_add(float* a, float* b, float* c, size_t N)
  {
      thrust::host_vector&lt;float&gt; h_a(a, a + N);
      thrust::host_vector&lt;float&gt; h_b(b, b + N);
  
      thrust::device_vector&lt;float&gt; d_a = h_a;
      thrust::device_vector&lt;float&gt; d_b = h_b;
      thrust::device_vector&lt;float&gt; d_c(N);
  
      vector_add(d_a, d_b, d_c);
  
      thrust::host_vector&lt;float&gt; h_c = d_c;
      for (int i = 0; i &lt; N; ++i)
      {
          c[i] = h_c[i];
      }
  }
  </code></pre>
  
  <p><code>lib.hpp</code>,</p>
  
  <pre><code>#pragma once
  
  #include &lt;cstddef&gt;
  
  void lib_vector_add(float*, float*, float*, size_t);
  </code></pre>
  
  <p><code>kernel.cuh</code> - this form results in a linker error. Uncomment the first <code>inline</code> to get a working code.</p>
  
  <pre><code>#pragma once
  
  #include &lt;thrust/device_vector.h&gt;
  #include &lt;cstddef&gt;
  
  // inline keyword avoids multiple definition errors, but produces warnings.
  // UNCOMMENT TO GET A WORKING EXECUTABLE.
  // inline
  __global__ void vector_add_kernel(
      const float *const a,
      const float *const b,
      float *const c,
      const size_t N)
  {
      int tid = threadIdx.x + blockIdx.x * blockDim.x;
  
      while (tid &lt; N)
      {
          c[tid] = a[tid] + b[tid];
          tid += blockDim.x * gridDim.x;
      }
  }
  
  // inline produces no warnings.
  inline
  void vector_add(
      const thrust::device_vector&lt;float&gt;&amp; d_a,
      const thrust::device_vector&lt;float&gt;&amp; d_b,
      thrust::device_vector&lt;float&gt;&amp; d_c)
  {
      const float *const a_ptr = thrust::raw_pointer_cast(d_a.data());
      const float *const b_ptr = thrust::raw_pointer_cast(d_b.data());
      float *const c_ptr = thrust::raw_pointer_cast(d_c.data());
  
      const size_t N = d_a.size();
  
      dim3 block(128);
      dim3 grid((N + 127) / 128);
  
      vector_add_kernel&lt;&lt;&lt;grid, block&gt;&gt;&gt;(a_ptr, b_ptr, c_ptr, N);
  }
  </code></pre>
  
  <p><code>Makefile</code></p>
  
  <pre><code>OBJS = main.o lib.o
  DEPS = kernel.cuh
  CU_ARCH = -gencode arch=compute_20,code=sm_20
  
  all: app
  
  app: $(OBJS)
      nvcc $(CU_ARCH) $(OBJS) -o app
  
  %.o: %.cu $(DEPS)
      nvcc $(CU_ARCH) -dc -I./ $&lt; -o $@
  
  clean:
      -rm *.o
  </code></pre>