title: Shared memory matrix multiplication kernel
body: <p>I am attempting to implement a shared memory based matrix multiplication kernel as outlined in the CUDA C Programming Guide. The following is the kernel:</p>
  
  <pre><code> __global__ void matrixMultiplyShared(float * A, float * B, float * C,
                       int ARows, int AColumns,
                       int BRows, int BColumns,
                       int CRows, int CColumns) {
       float * CSub = &amp;C[CColumns * 16 * blockIdx.y + 16 * blockIdx.x];
       float CValue = 0;
   for (int k = 0; k &lt; (AColumns / 16); ++k) {
           float * ASub =  &amp;A[AColumns * 16 * blockIdx.y + 16 * k];
           float * BSub = &amp;B[AColumns*16*k + 16*blockIdx.y];
           __shared__ float As[16][16];
           __shared__ float Bs[16][16];
           As[threadIdx.y][threadIdx.x] = ASub[threadIdx.y*AColumns+threadIdx.x];
           Bs[threadIdx.y][threadIdx.x] = BSub[threadIdx.y*AColumns+threadIdx.x];
           __syncthreads();
           for (int n = 0; n &lt; 16; ++n)
          CValue += As[threadIdx.y][n] * Bs[n][threadIdx.x];
           __syncthreads();
       }
       CSub[threadIdx.x*CColumns+threadIdx.y]=CValue;
   }
  </code></pre>
  
  <p>While the following is the call to the kernel:</p>
  
  <pre><code> dim3 dimBlock(16, 16, 1);
   dim3 dimGrid;
   dimGrid.x = (CColumns + dimBlock.x - 1)/dimBlock.x;
   dimGrid.y = (CRows + dimBlock.y - 1)/dimBlock.y;
   matrixMultiplyShared&lt;&lt;&lt;dimGrid , dimBlock&gt;&gt;&gt;(deviceA , deviceB , deviceC , ARows , AColumns, BRows ,BColumns , CRows , CColumns);
  </code></pre>
  
  <p>Unfortunately this seems to produce incorrect results.</p>
  
  <p>Any assistance/explanations would be greatly appreciated.</p>