title: Matrix vector product CUDA improve performance with tiling and shared memory
body: <p>Hello I'm working in a CUDA kernel about matrix vector product. I want to improve the performance with tiling and shared memory.
  The problem is that with this code the M Matrix or the N vector aren't loading right. </p>
  
  <p>Do you have any idea about how to Load a tile from M and N into the shared memory arrays??</p>
  
  <p>M is the matrix, N is the vector and P is the result of the matrix vector product</p>
  
  <pre><code>    __global__ void matrixMul( float* P, float* M, float* N, int Mw, int Nw)
  {
  int bx = blockIdx.x;     int by = blockIdx.y;
  int tx = threadIdx.x;    int ty = threadIdx.y;
  __shared__ float Ms[BLOCK_SIZE][BLOCK_SIZE];
  __shared__ float Ns[BLOCK_SIZE];
  
  // ===================================================================
  // Code segment 1
  // Determine the update values for the tile indices in the loop
  // ===================================================================
  
  int mBegin = Mw * BLOCK_SIZE * by;
  int mEnd   = mBegin + Mw - 1;
  int mStep  = BLOCK_SIZE;
  int nBegin = BLOCK_SIZE * bx;
  //int nStep  = BLOCK_SIZE*Nw;
  int nStep = 1;
  float Psub = 0.0f;
  
  // ===================================================================
  // Code segment 2
  // Do matrix-matrix multiplication inside a tile
  // ===================================================================
  
  for (int m = mBegin, n = nBegin; m &lt;= mEnd; m += mStep, n += nStep) {
  
      // Load a tile from M and N into the shared memory arrays
      Ms[ty][tx] = M[bx*mStep*Mw+m];
      Ns[ty] = N[by*nStep*Nw+n];
  
      // Synchronize the threads
      __syncthreads();
  
      // Multiply the two tiles together, each thread accumulating
      // the partial sum of a single dot product.
      for (int i = 0; i &lt; BLOCK_SIZE; i++) {
          Psub += Ms[i][tx] * Ns[i];
      }
  
      // Synchronize again.
      __syncthreads();
  }
  
  // ===================================================================
  // Code segment 3
  // Store the data back to global memory
  // ===================================================================
  
  int p = Nw * BLOCK_SIZE * by + BLOCK_SIZE * bx;
  P[p + nStep] = Psub;
  }
  </code></pre>