title: CUDA __device__ Unresolved extern function
body: <p>I am trying to understand how to decouple CUDA <code>__device__</code> codes in separate header files.</p>
  <p>I have three files.</p>
  <p><strong>File: 1: int2.cuh</strong></p>
  <pre><code>#ifndef INT2_H_
  #define INT2_H_
  
  #include &quot;cuda.h&quot;
  #include &quot;cuda_runtime.h&quot;
  #include &quot;device_launch_parameters.h&quot;
  
  __global__ void kernel();
  __device__ int k2(int k);
  
  int launchKernel(int dim);
  
  #endif /* INT2_H_ */
  </code></pre>
  <p><strong>File 2: int2.cu</strong></p>
  <pre><code>#include &quot;int2.cuh&quot;
  #include &quot;cstdio&quot;
  
  __global__ void kernel() {
      int tid = threadIdx.x;
      printf(&quot;%d\n&quot;, k2(tid));
  }
  
  __device__ int k2(int i) {
      return i * i;
  }
  
  int launchKernel(int dim) {
      kernel&lt;&lt;&lt;1, dim&gt;&gt;&gt;();
      cudaDeviceReset();
      return 0;
  }
  </code></pre>
  <p><strong>File 3: CUDASample.cu</strong></p>
  <pre><code>include &lt;stdio.h&gt;
  #include &lt;stdlib.h&gt;
  #include &quot;int2.cuh&quot;
  #include &quot;iostream&quot;
  
  using namespace std;
  
  static const int WORK_SIZE = 256;
  
  __global__ void sampleCuda() {
      int tid = threadIdx.x;
  //    printf(&quot;%d\n&quot;, k2(tid)); //Can not call k2
      printf(&quot;%d\n&quot;, tid * tid);
  }
  
  int main(void) {
  
      int var;
      var = launchKernel(16);
  
      kernel&lt;&lt;&lt;1, 16&gt;&gt;&gt;();
      cudaDeviceReset();
  
      sampleCuda&lt;&lt;&lt;1, 16&gt;&gt;&gt;();
      cudaDeviceReset();
  
      return 0;
  }
  </code></pre>
  <p>The code works fine. I can call the <code>sampleCuda()</code> kernel (in same file), call the C function <code>launchKernel()</code> (in other file), and call <code>kernel()</code> directly (in other file).</p>
  <p>However, I get the following error when calling the <code>__device__</code> function from the <code>sampleCuda()</code> kernel. The same function is callable in <code>kernel()</code>.</p>
  <pre><code>10:58:11 **** Incremental Build of configuration Debug for project CUDASample ****
  make all 
  Building file: ../src/CUDASample.cu
  Invoking: NVCC Compiler
  /Developer/NVIDIA/CUDA-6.5/bin/nvcc -G -g -O0 -gencode arch=compute_20,code=sm_20  -odir &quot;src&quot; -M -o &quot;src/CUDASample.d&quot; &quot;../src/CUDASample.cu&quot;
  /Developer/NVIDIA/CUDA-6.5/bin/nvcc -G -g -O0 --compile --relocatable-device-code=false -gencode arch=compute_20,code=compute_20 -gencode arch=compute_20,code=sm_20  -x cu -o  &quot;src/CUDASample.o&quot; &quot;../src/CUDASample.cu&quot;
  ../src/CUDASample.cu(18): warning: variable &quot;var&quot; was set but never used
  
  ../src/CUDASample.cu(8): warning: variable &quot;WORK_SIZE&quot; was declared but never referenced
  
  ../src/CUDASample.cu(18): warning: variable &quot;var&quot; was set but never used
  
  ../src/CUDASample.cu(8): warning: variable &quot;WORK_SIZE&quot; was declared but never referenced
  
  ptxas fatal   : Unresolved extern function '_Z2k2i'
  make: *** [src/CUDASample.o] Error 255
  
  10:58:14 Build Finished (took 2s.388ms)
  </code></pre>
  <p>How do I call the <code>__device__</code> function from the <code>sampleCuda()</code> kernel ?</p>