title: CUDA Matrix multiplication breaks for large matrices
body: <p>I have the following matrix multiplication code, implemented using CUDA 3.2 and VS 2008. I am running on Windows server 2008 r2 enterprise. I am running a Nvidia GTX 480. The following code works fine with values of "Width" (Matrix width) up to about 2500 or so.</p>
  
  <pre><code>int size = Width*Width*sizeof(float);
  float* Md, *Nd, *Pd;
  cudaError_t err = cudaSuccess;
  
  //Allocate Device Memory for M, N and P
  err = cudaMalloc((void**)&amp;Md, size);
  err = cudaMalloc((void**)&amp;Nd, size);
  err = cudaMalloc((void**)&amp;Pd, size);
  
  //Copy Matrix from Host Memory to Device Memory
  err = cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice);
  err = cudaMemcpy(Nd, N, size, cudaMemcpyHostToDevice);
  
  //Setup the execution configuration
  dim3 dimBlock(TileWidth, TileWidth, 1);
  dim3 dimGrid(ceil((float)(Width)/TileWidth), ceil((float)(Width)/TileWidth), 1);
  
  MatrixMultiplicationMultiBlock_Kernel&lt;&lt;&lt;dimGrid, dimBlock&gt;&gt;&gt;(Md, Nd, Pd, Width);
  
  err = cudaMemcpy(P, Pd, size, cudaMemcpyDeviceToHost);
  
  //Free Device Memory
  cudaFree(Md);
  cudaFree(Nd);
  cudaFree(Pd);
  </code></pre>
  
  <p>When I set the "Width" to 3000 or greater, I get the following error after a black screen:
  <img src="https://i.stack.imgur.com/zGhB7.png" alt="screenshot"></p>
  
  <p>I looked online and I saw that some people has this issue because the watchdog was killing the kernel after it hangs for more than 5 seconds. I tried editing the "TdrDelay" in the registry and this delayed the time before the black screen and same error appeared. So I concluded this was not my issue.</p>
  
  <p>I debugged into my code and found this line to be the culprit:</p>
  
  <pre><code>err = cudaMemcpy(P, Pd, size, cudaMemcpyDeviceToHost);
  </code></pre>
  
  <p>This is what I use to return my result set from the device after my matrix multiplication kernel function is called. Everything up until this point seems to run fine. I believe I am allocating memory correctly and cannot figure out why this is happening. I thought maybe I didn't have enough memory on my card for this but then shouldn't cudaMalloc have returned an error? (I confirmed it didn't while debugging).</p>
  
  <p>Any ideas/assistance would be greatly appreciated!... Thanks a lot guys!!</p>
  
  <p>Kernel code:</p>
  
  <pre><code>//Matrix Multiplication Kernel - Multi-Block Implementation
  __global__ void MatrixMultiplicationMultiBlock_Kernel (float* Md, float* Nd, float* Pd, int Width) 
  {
  int TileWidth = blockDim.x;
  
  //Get row and column from block and thread ids
  int Row = (TileWidth*blockIdx.y) + threadIdx.y;
  int Column = (TileWidth*blockIdx.x) + threadIdx.x;
  
  //Pvalue store the Pd element that is computed by the thread
  float Pvalue = 0;
  
  for (int i = 0; i &lt; Width; ++i)
  {
      float Mdelement = Md[Row * Width + i];
      float Ndelement = Nd[i * Width + Column];
      Pvalue += Mdelement * Ndelement;
  }
  
  //Write the matrix to device memory each thread writes one element
  Pd[Row * Width + Column] = Pvalue;
  }
  </code></pre>
  
  <p>I also have this other function that uses shared memory, and it also gives the same error:</p>
  
  <p>Call:</p>
  
  <pre><code>            MatrixMultiplicationSharedMemory_Kernel&lt;&lt;&lt;dimGrid, dimBlock, sizeof(float)*TileWidth*TileWidth*2&gt;&gt;&gt;(Md, Nd, Pd, Width);
  </code></pre>
  
  <p>Kernel code:</p>
  
  <pre><code> //Matrix Multiplication Kernel - Shared Memory Implementation
   __global__ void MatrixMultiplicationSharedMemory_Kernel (float* Md, float* Nd, float* Pd, int Width) 
   {
  int TileWidth = blockDim.x;
  
  //Initialize shared memory
  extern __shared__ float sharedArrays[];
  float* Mds = (float*) &amp;sharedArrays;
  float* Nds = (float*) &amp;Mds[TileWidth*TileWidth];
  
  int tx = threadIdx.x;
  int ty = threadIdx.y;
  
  //Get row and column from block and thread ids
  int Row = (TileWidth*blockIdx.y) + ty;
  int Column = (TileWidth*blockIdx.x) + tx;
  float Pvalue = 0;
  
  //For each tile, load the element into shared memory
  for( int i = 0; i &lt; ceil((float)Width/TileWidth); ++i)
  {
      Mds[ty*TileWidth+tx] = Md[Row*Width + (i*TileWidth + tx)];
      Nds[ty*TileWidth+tx] = Nd[(ty + (i * TileWidth))*Width + Column]; 
  
      __syncthreads();
  
      for( int j = 0; j &lt; TileWidth; ++j)
      {
          Pvalue += Mds[ty*TileWidth+j] * Nds[j*TileWidth+tx];
      }
  
      __syncthreads();
  }
  
  //Write the matrix to device memory each thread writes one element
  Pd[Row * Width + Column] = Pvalue;
  }
  </code></pre>