title: Limiting register usage in CUDA: __launch_bounds__ vs maxrregcount
body: <p>From the <a href="http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#ixzz4kk6SRKqY" rel="noreferrer">NVIDIA CUDA C Programming Guide</a>: </p>
  
  <blockquote>
    <p>Register usage can be controlled using the <code>maxrregcount</code> compiler
    option or launch bounds as described in Launch Bounds.</p>
  </blockquote>
  
  <p>From my understanding (and correct me if I'm wrong), while <code>-maxrregcount</code> limits the number of registers the entire <code>.cu</code> file may use, the <code>__launch_bounds__</code> qualifier defines the <code>maxThreadsPerBlock</code> and <code>minBlocksPerMultiprocessor</code> for each <code>__global__</code> kernel. These two accomplish the same task, but in two different ways.</p>
  
  <p>My usage requires me to have <code>40</code> registers per thread to maximize the performance. Thus, I can use <code>-maxrregcount 40</code>. I can also force <code>40</code> registers by using <code>__launch_bounds__(256, 6)</code> but this causes load &amp; store register spills. </p>
  
  <p>What is the difference between the two to cause these register spills?</p>