title: How to pass an array of structs to CUDA kernel using constant memory
body: <p>My kernel needs a list/array of Configuration as an input parameter. I even have a list/array of such lists/arrays, one of them is to pass to the kernel. These Configuration are prepared on the host and do not change. So this would be a perfect use for constant memory. But honestly, I do not really get how to do it.</p>
  
  <p>I try to give my idea in the code draft below. Basically, I see two ways how to define/pass the lists:</p>
  
  <ul>
  <li>Define them as arrays with fixed lenghts and pass them by-value to the kernel</li>
  <li>Define them as pointers and just pass a pointer to the kernel (must be copied to device first, of course)</li>
  </ul>
  
  <p>Which method should I take and how should I modify the code below to make sure, constant memory is used?</p>
  
  <p>I expect each list to have typically a size less than 200-300 Bytes. If I would make all lists of the same size, I would maybe go for a size of 512 Bytes or 1 kB.</p>
  
  <pre><code>class Configuration{
    // some constants
  }
  
  // We need a list of lists Configurations, these could be implemented either as...
  Configuration a[10][100]; // fixed-length array or...
  Configuration ** b; // as a dynamic array to pointers of arrays
  
  // Parameter will take an array of Configuration, either as a pointer or directly as an array
  __global__ kernel(Configuration * config){
  
  }
  
  // According to the above example, we use the pointer-version. Could also be a call directly using a[i]
  kernel&lt;&lt;...&gt;&gt;(b[i], lengthOfB[i]);
  </code></pre>