title: Shared Memory Bank Conflicts in CUDA: How memory is aligned to banks
body: <p>As far as my understanding goes, shared memory is divided into banks and accesses by multiple threads to a single data element within the same bank will cause a conflict (or broadcast).</p>
  
  <p>At the moment I allocate a fairly large array which conceptually represents several pairs of two matrices:</p>
  
  <pre><code>__shared__ float A[34*N]
  </code></pre>
  
  <p>Where <code>N</code> is the number of pairs and the first 16 floats of a pair are one matrix and the following 18 floats are the second.</p>
  
  <p>The thing is, access to the first matrix is conflict free but access to the second one has conflicts.  These conflicts are unavoidable, however, my thinking is that because the second matrix is 18 all future matrices will be misaligned to the banks and therefore more conflicts than necessary will occur.</p>
  
  <p>Is this true, if so how can I avoid it?</p>
  
  <p>Everytime I allocate shared memory, does it start at a new bank? So potentially could I do</p>
  
  <pre><code>__shared__ Apair1[34]
  __shared__ Apair2[34]
  ...
  </code></pre>
  
  <p>Any ideas?</p>
  
  <p>Thanks</p>