title: Row-wise/column-wise operations on matrices with CUDA
body: <p>I'm relatively new to CUDA programming. I have understood the programming model and have already written few basic kernels. I know how to apply a kernel to each element of a matrix (stored as 1D array), but now I'm trying to figure out <strong>how to apply the same operation to the same row/column of the input matrix</strong>. </p>
  
  <p>Let's say I have a <code>MxN</code> matrix and a vector of length <code>N</code>. I would like to sum (but it can be any other math operation) the vector to each row of the matrix. 
  The serial code of such operation is: </p>
  
  <pre class="lang-cpp prettyprint-override"><code>for (int c = 0; c &lt; columns; c++) 
  {
      for (int r = 0; r &lt; rows; r++)
      {
          M[r * rows + c] += V[c];
      }
  }
  </code></pre>
  
  <p>Now the CUDA code for doing this operation should be quite straightforward: I should spawn as many cuda threads as the elements and apply this kernel:</p>
  
  <pre class="lang-cpp prettyprint-override"><code>__global__ void kernel(const unsigned int size, float* matrix, const float* vector)
  {
      // get the current element index for the thread
      unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
      if (idx &lt; size)
      {
          // sum the current element with the 
          matrix[idx] += vector[threadIdx.x];
      }
  }
  </code></pre>
  
  <p>It runs but the result is not correct. Actually, it's correct if I transpose the matrix after the kernel completes its work. Unfortunately, I have no clue why it works in this way. Could you help me to figure out this problem? Thanks in advance.</p>
  
  <p><em>EDIT #1</em></p>
  
  <p>I launch the kernel using:</p>
  
  <pre class="lang-cpp prettyprint-override"><code>int block_size = 64;
  int grid_size = (M * N + block_size - 1) / block_size;
  kernel&lt;&lt;&lt;grid_size, block_size&gt;&gt;&gt;(M * N, matrix, vector);
  </code></pre>
  
  <p><em>EDIT #2</em></p>
  
  <p>I solved the problem by fixing the CPU code as suggested by @RobertCrovella:</p>
  
  <pre class="lang-cpp prettyprint-override"><code>M[r * columns + c] += V[c];
  </code></pre>
  
  <p>It should match the outer <code>for</code>, that is, over the columns.</p>