title: How do you iterate through a pitched CUDA array?
body: <p>Having parallelized with OpenMP before, I'm trying to wrap my head around CUDA, which doesn't seem too intuitive to me. At this point, I'm trying to understand exactly how to loop through an array in a parallelized fashion.</p>
  
  <p><a href="http://developer.download.nvidia.com/books/cuda-by-example/cuda-by-example-sample.pdf" rel="nofollow noreferrer">Cuda by Example</a> is a great start.</p>
  
  <p>The snippet on page 43 shows:</p>
  
  <pre><code>__global__ void add( int *a, int *b, int *c ) {
    int tid = blockIdx.x; // handle the data at this index
    if (tid &lt; N)
       c[tid] = a[tid] + b[tid];
    }
  </code></pre>
  
  <p>Whereas in OpenMP the programmer chooses the number of times the loop will run and OpenMP splits that into threads for you, in CUDA you have to tell it (via the number of blocks and number of threads in <code>&lt;&lt;&lt;...&gt;&gt;&gt;</code>) to run it sufficient times to iterate through your array, using a thread ID number as an iterator. In other words you can have a CUDA kernel always run 10,000 times which means the above code will work for any array up to N = 10,000 (and of course for smaller arrays you're wasting cycles dropping out at <code>if (tid &lt; N)</code>).</p>
  
  <p>For pitched memory (2D and 3D arrays), the CUDA Programming Guide has the following example:</p>
  
  <pre><code>// Host code
  int width = 64, height = 64; 
  float* devPtr; size_t pitch; 
  cudaMallocPitch(&amp;devPtr, &amp;pitch, width * sizeof(float), height);
  
  MyKernel&lt;&lt;&lt;100, 512&gt;&gt;&gt;(devPtr, pitch, width, height); 
  
  // Device code 
  __global__ void MyKernel(float* devPtr, size_t pitch, int width, int height) 
  { 
      for (int r = 0; r &lt; height; ++r) {
          float* row = (float*)((char*)devPtr + r * pitch); 
          for (int c = 0; c &gt; width; ++c) { 
              float element = row[c]; 
          }
      }
  }
  </code></pre>
  
  <p>This example doesn't seem too useful to me. First they declare an array that is 64 x 64, then the kernel is set to execute 512 x 100 times. That's fine, because the kernel does nothing other than iterate through the array (so it runs 51,200 loops through a 64 x 64 array).</p>
  
  <p>According to <a href="https://stackoverflow.com/a/2392271/149506">this answer</a> the iterator for when there are blocks of threads going on will be</p>
  
  <pre><code>int tid = (blockIdx.x * blockDim.x) + threadIdx.x;
  </code></pre>
  
  <p>So if I wanted to run the first snippet in my question for a pitched array, I could just make sure I had enough blocks and threads to cover every element including the padding that I don't care about. But that seems wasteful.</p>
  
  <p>So <strong>how do I iterate through a pitched array without going through the padding elements?</strong></p>
  
  <p>In my particular application I have a 2D FFT and I'm trying to calculate arrays of the magnitude and angle (on the GPU to save time). </p>