title: CUDA C++ Templating of Kernel Parameter
body: <p>I'm trying to templatize a CUDA kernel based on a boolean variable (as shown here: <a href="https://stackoverflow.com/questions/6179295/if-statement-inside-a-cuda-kernel">Should I unify two similar kernels with an &#39;if&#39; statement, risking performance loss?</a>), but I keep getting a compiler error that says my function is not a template. I think that I'm just missing something obvious so it's pretty frustrating.</p>
  
  <p>The following does NOT work:</p>
  
  <p>util.cuh</p>
  
  <pre><code>#include "kernels.cuh"
  //Utility functions
  </code></pre>
  
  <p>kernels.cuh</p>
  
  <pre><code>    #ifndef KERNELS
      #define KERNELS
      template&lt;bool approx&gt;
      __global__ void kernel(...params...);
      #endif
  </code></pre>
  
  <p>kernels.cu</p>
  
  <pre><code>template&lt;bool approx&gt;
  __global__ void kernel(...params...)
  {
      if(approx)
      {
         //Approximate calculation
      }
      else
      {
        //Exact calculation
      }
  }
  
  template __global__ void kernel&lt;false&gt;(...params...); //Error occurs here
  </code></pre>
  
  <p>main.cu</p>
  
  <pre><code>#include "kernels.cuh"
  kernel&lt;false&gt;&lt;&lt;&lt;dimGrid,dimBlock&gt;&gt;&gt;(...params...);
  </code></pre>
  
  <p>The following DOES work:</p>
  
  <p>util.cuh</p>
  
  <pre><code>#include "kernels.cuh"
  //Utility functions
  </code></pre>
  
  <p>kernels.cuh</p>
  
  <pre><code>#ifndef KERNELS
  #define KERNELS
  template&lt;bool approx&gt;
  __global__ void kernel(...params...);
  template&lt;bool approx&gt;
  __global__ void kernel(...params...)
  {
      if(approx)
      {
         //Approximate calculation
      }
      else
      {
        //Exact calculation
      }
  }
  #endif
  </code></pre>
  
  <p>main.cu</p>
  
  <pre><code>#include "kernels.cuh"
  kernel&lt;false&gt;&lt;&lt;&lt;dimGrid,dimBlock&gt;&gt;&gt;(...params...);
  </code></pre>
  
  <p>If I throw in the </p>
  
  <pre><code>template __global__ void kernel&lt;false&gt;(...params...);
  </code></pre>
  
  <p>line at the end of kernels.cuh it also works.</p>
  
  <p>I get the following errors (both referring to the marked line above):</p>
  
  <pre><code>kernel is not a template
  invalid explicit instantiation declaration
  </code></pre>
  
  <p>If it makes a difference I compile all of my .cu files in one line, like:</p>
  
  <pre><code>nvcc -O3 -arch=sm_21 -I. main.cu kernels.cu -o program
  </code></pre>