title: Namespaces as template parameters in CUDA
body: <p>In C++, it is impossible to pass a <code>namespace</code> as some sort of parameter (by means of templates or actual function parameters) to a class or function is not possible. The same applies to CUDA (at least, to my knowledge). Some reasons are explained in this question: <a href="https://stackoverflow.com/questions/12905951/why-cant-namespaces-be-template-parameters">Why can&#39;t namespaces be template parameters?</a></p>
  
  <p>Here is an example of a use case:</p>
  
  <pre><code>namespace experiment1
  {
      int repetitions() { return 2; }
      void setup() { ... }
      void f() { ... }
      void teardown() { ... }
  }
  
  namespace experiment2
  {
      int repetitions() { return 4; }
      void setup() { ... }
      void f() { ... }
      void teardown() { ... }
  }
  
  // Beware, this is invalid C++ and invalid CUDA
  template&lt;namespace NS&gt;
  void do_test()
  {
      // Do something with NS::repetitions(), NS::setup(), ...
  } 
  </code></pre>
  
  <p>One of the reasons why this is not valid in C++ is that there is nothing in this approach that you could not do with classes. You could indeed turn every namespace into a class and the functions into member functions, then pass the class as template parameter to the <code>do_test</code> function, or an instance of it as parameter to the same function (possibly using static functions in the former case or virtual functions in the latter case).</p>
  
  <p>I agree with this. However, in the specific case of CUDA, there is something you could do with namespaces, but not with classes. Imagine <code>f</code> is a kernel, i.e. a <code>__global__</code> function, and that <code>setup</code> or another function is used to specify, e.g. the size of the shared memory to be allocated for the kernel. Kernels cannot be members of classes (see the answer to this question: <a href="https://stackoverflow.com/questions/31120692/can-cuda-kernels-be-virtual-functions">Can CUDA kernels be virtual functions?</a>). However, you can enclose it with the other functions related to the same experiment in the same <code>namespace</code>.</p>
  
  <p>Consider the case shown in the code above: <code>do_test</code> is the function which sets up the timers, prepares some input, checks the output, measures the time and does some other operations. Every experiment is a set of a few functions, with the same name and same interface, one of which is a kernel. You want <code>do_test</code> to be generic enough to handle all these experiments. And you want the code for each experiment to be self-contained in some form of encapsulation like namespace, struct, class,...</p>
  
  <p>Can this problem be solved?</p>
  
  <hr>
  
  <p>As requested by talonmies (many thanks for your comment, by the way), I will make the question more concrete.</p>
  
  <p>I have several very simple kernels which perform similar operations. They load values form one large array, they apply a stencil operation on them and they write the result onto an output array (different from the input array). By stencil operation, I mean an operation performed by thread <code>idx</code> on input values <code>idx</code>, and its neighboring values (say from <code>idx-3</code> to <code>idx+3</code>). The simplest of these kernels performs just a copy from input to output: every thread reads <code>input[idx]</code> and writes <code>output[idx]</code>. Another example is a difference stencil which performs <code>output[idx] = input[idx+1] - input[idx-1]</code>. (I'm leaving apart some details, but you get the idea.)</p>
  
  <p>I want to benchmark these kernels so to derive a performance model. For each kernel I also need a host function which is able to check the result. I also have in each case another kernel which performs the same operation in a slightly different way through an optimization, but equivalent from the point of view of the result. Finally, I have a host function which prints the name of the kernel. Here is an summary in code:</p>
  
  <pre><code>namespace copy
  {
      std::string name() { return "copy"; }
      __global__ void kernel(const float* input, float* output, int size);
      __global__ void kernelOptimized(const float* input, float* output, int size);
      bool check(const float* input, const float* output);
  }
  
  namespace difference
  {
      std::string name() { return "difference"; }
      __global__ void kernel(const float* input, float* output, int size);
      __global__ void kernelOptimized(const float* input, float* output, int size);
      bool check(const float* input, const float* output);
  }
  </code></pre>
  
  <p>I have a function <code>do_test</code>, that I parametrized to be generic:</p>
  
  <pre><code>typedef bool NameFunction(const float* input, const float* output);
  typedef bool CheckFunction(const float* input, const float* output);
  typedef void KernelFunction(const float* input, float* output, int size);
  
  void do_test(NameFunction name, KernelFunction kernel1, KernelFunction kernel2, CheckFunction check)
  {
      // Set up input and output array
      // Set up CUDA events
      // Warm up kernels
      // Run kernels
      // Check results
      // Measure time
      // Do standard output
  }
  
  int main()
  {
      do_test&lt;copy::name, copy::kernel, copy::kernelOptimized, copy::check&gt;()
      do_test&lt;difference::name, difference::kernel, difference::kernelOptimized, difference::check&gt;()
  }
  </code></pre>
  
  <p>Now, of course this way is already quite good. However, if I introduce one more function that every experiment must provide, I will need to modify all these lines where I call <code>do_test</code>. I would prefer passing this namespace or some sort of object containing these functions.</p>