title: Smart design for large kernel with different inputs that only changes one line of code
body: <p>I am designing some kernels that I would like to have 2 ways of calling: Once with standard <code>float *</code> device as input (for writing), and another with <code>cudaSurfaceObject_t</code> as input (for writing). The kernel itself is long (>200 lines) and ultimately, I only need the last line to be different. In one case you have standard <code>out[idx]=val</code> type of assignment, while in the other one a <code>surf3Dwrite()</code> type. The rest of the kernel is identical. </p>
  
  <p>Something like </p>
  
  <pre><code>__global__ kernel(float * out , ....)
  {
  
  // 200 lines of math
  
  // only difference, aside from input argument
  idx=....
  out[idx]=a;
  }
  </code></pre>
  
  <p>vs</p>
  
  <pre><code>__global__ kernel(cudaSurfaceObject_t *  out, ...)
  {
  
  // 200 lines of math
  
  // only difference, aside from input argument
    surf3Dwrite(&amp;out,val,x,y,z);
  }
  </code></pre>
  
  <p>What is the smart way of coding this, without copy pasting the entire kernel and renaming it? I checked Templating, but (if I am not wrong) its for types only, one can not just have a completely different line of code when the type is different in a template. CUDA kernels don't seem to be able to be overloaded either. </p>