title: Efficiently pass parameters to CUDA Kernel
body: <p>I'm currently refactoring my CUDA code and stumbled across a problem: Passing multiple device pointers to the kernel can become quite ugly and, therefore, I thought of grouping them together inside a struct of arrays like so:</p>
  <pre><code>struct Data {
      float* a;
      float* b;
      float* c;
      // ...
  };
  
  __global__ void foo(Data* data, size_t n) {
      size_t const id = blockIdx.x * blockDim.x + threadIdx.x;
      if (id &lt; n) {
          data-&gt;c[id] = data-&gt;a[id] + data-&gt;b[id];
      }
  }
  </code></pre>
  <p>Now consider this host code:</p>
  <pre><code>struct Data host_data;
  struct Data device_data;
  
  // allocate host memory
  size_t const kSize = 10;
  host_data.a = (float*)malloc(sizeof(float) * kSize);
  host_data.b = (float*)malloc(sizeof(float) * kSize);
  host_data.c = (float*)malloc(sizeof(float) * kSize);
  
  // allocate device memory
  cudaMalloc((void **)&amp;device_data.a, sizeof(float) * kSize);
  cudaMalloc((void **)&amp;device_data.b, sizeof(float) * kSize);
  cudaMalloc((void **)&amp;device_data.c, sizeof(float) * kSize);
  
  // initialize some host data
  for (size_t i = 0; i &lt; kSize; ++i) {
    host_data.a[i] = 10.f;
    host_data.b[i] = 10.f;
  }
  
  // copy host data to device data
  cudaMemcpy(device_data.a, host_data.a, sizeof(float) * kSize, cudaMemcpyHostToDevice);
  cudaMemcpy(device_data.b, host_data.b, sizeof(float) * kSize, cudaMemcpyHostToDevice);
  cudaMemcpy(device_data.c, host_data.c, sizeof(float) * kSize, cudaMemcpyHostToDevice);
  
  // make sure to allocate the device_ptr struct itself on the device memory
  struct Data* device_ptr;
  cudaMalloc((void**)&amp;device_ptr, sizeof(*device_ptr));
  
  // bind local device pointers
  cudaMemcpy(&amp;(device_ptr-&gt;a), &amp;device_data.a, sizeof(device_ptr-&gt;a), cudaMemcpyHostToDevice);
  cudaMemcpy(&amp;(device_ptr-&gt;b), &amp;device_data.b, sizeof(device_ptr-&gt;b), cudaMemcpyHostToDevice);
  cudaMemcpy(&amp;(device_ptr-&gt;c), &amp;device_data.c, sizeof(device_ptr-&gt;c), cudaMemcpyHostToDevice);
  
  // launch kernel
  size_t const block_size = 512;
  size_t const grid_size = ceil((float)kSize/block_size);
  simple &lt;&lt;&lt; grid_size, block_size &gt;&gt;&gt; (device_ptr, kSize);
  
  // wait for kernel to finish
  cudaDeviceSynchronize();
  
  // copy device data back to host data
  cudaMemcpy(host_data.a, device_data.a, sizeof(float) * kSize, cudaMemcpyDeviceToHost);
  cudaMemcpy(host_data.b, device_data.b, sizeof(float) * kSize, cudaMemcpyDeviceToHost);
  cudaMemcpy(host_data.c, device_data.c, sizeof(float) * kSize, cudaMemcpyDeviceToHost);
  </code></pre>
  <p>This of course works, but it seems counterintuitive, as sharing data structures between the host and device becomes extremely tedious. Is there a way to avoid / reduce the amount of <code>cudaMemcpy</code> and <code>cudaMalloc</code> calls somehow?</p>
  <p>FYI: I'm aware that passing a copy of the <code>Data</code> struct to the kernel is possible. However, this is not very efficient for large structs.</p>