title: Un-coalesced global memory access caused by indirect access in CUDA
body: <p>My CUDA programm is suffering from un-coalesced global memory access. Although the idx-th thread only deal with the [idx]-th cell in an array, there are many indirect memory accesses as shown below. </p>
  
  <pre><code>int idx=blockDim.x*blockIdx.x+threadIdx.x;
  
  .... = FF[m_front[m_fside[idx]]];
  </code></pre>
  
  <p>For m_fisde[idx], we have coalesced accesses, but what we actually need is FF[m_front[m_fside[idx]]]. There is a two-level indirect access.</p>
  
  <p>I tried to find some patterns of the data in m_front or m_fsied in order to make this to be a direct sequential access, but found out that they are almost 'random'. </p>
  
  <p>Is there a possible way to tackle this?</p>