title: Amount of local memory per CUDA thread
body: <p>I read in NVIDIA documentation (<a href="http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#features-and-technical-specifications" rel="nofollow">http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#features-and-technical-specifications</a>, table #12) that the amount of local memory per thread is 512 Ko for my GPU (GTX 580, compute capability 2.0).</p>
  
  <p>I tried unsuccessfully to check this limit on Linux with CUDA 6.5.</p>
  
  <p>Here is the code I used (its only purpose is to test local memory limit, it doesn't make any usefull computation):</p>
  
  <pre><code>#include &lt;iostream&gt;
  #include &lt;stdio.h&gt;
  
  #define MEMSIZE 65000  // 65000 -&gt; out of memory, 60000 -&gt; ok
  
  inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=false)
  {
      if (code != cudaSuccess) 
      {
          fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
          if( abort )
              exit(code);
      }
  }
  
  inline void gpuCheckKernelExecutionError( const char *file, int line)
  {
      gpuAssert( cudaPeekAtLastError(), file, line);
      gpuAssert( cudaDeviceSynchronize(), file, line);    
  }
  
  
  __global__ void kernel_test_private(char *output)
  {
      int c = blockIdx.x*blockDim.x + threadIdx.x; // absolute col
      int r = blockIdx.y*blockDim.y + threadIdx.y; // absolute row
  
      char tmp[MEMSIZE];
      for( int i = 0; i &lt; MEMSIZE; i++)
          tmp[i] = 4*r + c; // dummy computation in local mem
      for( int i = 0; i &lt; MEMSIZE; i++)
          output[i] = tmp[i];
  }
  
  int main( void)
  {
      printf( "MEMSIZE=%d bytes.\n", MEMSIZE);
  
      // allocate memory
      char output[MEMSIZE];
      char *gpuOutput;
      cudaMalloc( (void**) &amp;gpuOutput, MEMSIZE);
  
      // run kernel
      dim3 dimBlock( 1, 1);
      dim3 dimGrid( 1, 1);
      kernel_test_private&lt;&lt;&lt;dimGrid, dimBlock&gt;&gt;&gt;(gpuOutput);
      gpuCheckKernelExecutionError( __FILE__, __LINE__);
  
      // transfer data from GPU memory to CPU memory
      cudaMemcpy( output, gpuOutput, MEMSIZE, cudaMemcpyDeviceToHost);
  
      // release resources
      cudaFree(gpuOutput);
      cudaDeviceReset();
  
      return 0;
  }
  </code></pre>
  
  <p>And the compilation command line:</p>
  
  <pre><code>nvcc -o cuda_test_private_memory -Xptxas -v -O2 --compiler-options -Wall cuda_test_private_memory.cu
  </code></pre>
  
  <p>The compilation is ok, and reports:</p>
  
  <pre><code>ptxas info    : 0 bytes gmem
  ptxas info    : Compiling entry function '_Z19kernel_test_privatePc' for 'sm_20'
  ptxas info    : Function properties for _Z19kernel_test_privatePc
      65000 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
  ptxas info    : Used 21 registers, 40 bytes cmem[0]
  </code></pre>
  
  <p>I got an "out of memory" error at runtime on the GTX 580 when I reached 65000 bytes per thread. Here is the exact output of the program in the console:</p>
  
  <pre><code>MEMSIZE=65000 bytes.
  GPUassert: out of memory cuda_test_private_memory.cu 48
  </code></pre>
  
  <p>I also did a test with a GTX 770 GPU (on Linux with CUDA 6.5). It ran without error for MEMSIZE=200000, but the "out of memory error" occurred at runtime for MEMSIZE=250000.</p>
  
  <p>How to explain this behavior ? Am I doing something wrong ? </p>