title: CUDA shared memory bank conflicts report higher
body: <p>I've been working on optimizing some code and ran into an issue with the shared memory bank conflict report with the CUDA Nsight performance analysis.  I was able to reduce it to a very simple piece of code that Nsight reports as having a bank conflict, when it doesn't seem there should be one.  Below is the kernel:</p>
  
  <pre><code>__global__ void conflict() {
      __shared__ double values[33];
      values[threadIdx.x] = threadIdx.x;
      values[threadIdx.x+1] = threadIdx.x;
  }
  </code></pre>
  
  <p>And the main function to call it:</p>
  
  <pre><code>int main() {
      conflict&lt;&lt;&lt;1,32&gt;&gt;&gt;();
  }
  </code></pre>
  
  <p>Note that I am using a single warp to really reduce this to the bare minimum.  When I run the code, Nsight says there is 1 bank conflict, but according to everything I have read, there should not be any.  For each access to the shared memory array, each thread is accessing consecutive values, each belonging to separate banks.</p>
  
  <p>Has anyone else experienced issues with the reporting of Nsight or am I just missing something with the functioning of bank conflicts?  I would appreciate any feedback!</p>
  
  <p>Btw, I am running the following setup:</p>
  
  <ul>
  <li>Windows 8</li>
  <li>GTX 770</li>
  <li>Visual Studio Community 2013</li>
  <li>CUDA 7</li>
  <li>Nsight Visual studio edition Version 4.5</li>
  </ul>