title: How to properly coalesce reads from global memory into shared memory with elements of type short or char (assuming one thread per element)?
body: <p>I have a questions about coalesced global memory loads in CUDA. Currently I need to be able to execute on a CUDA device with compute capability CUDA 1.1 or 1.3.</p>
  
  <p>I am writing a CUDA kernel function which reads an array of type T from global memory into shared memory, does some computation, and then will write out an array of type T back to global memory. I am using the shared memory because the computation for each output element actually depends not only on the corresponding input element, but also on the nearby input elements. I only want to load each input element once, hence I want to cache the input elements in shared memory.</p>
  
  <p>My plan is to have each thread read one element into shared memory, then __syncthreads() before beginning the computation. In this scenario, each thread loads, computes, and stores one element (although the computation depends on elements loaded into shared memory by other threads).</p>
  
  <p>For this question I want to focus on the read from global memory into shared memory.</p>
  
  <p>Assuming that there are N elements in the array, I have configured CUDA to execute a total of N threads. For the case where sizeof(T) == 4, this should coalesce nicely according to my understanding of CUDA, since thread K will read word K (where K is the thread index).</p>
  
  <p>However, in the case where <code>sizeof(T) &lt; 4</code>, for example if T=unsigned char or if T=short, then I think there may be a problem. In this case, my (naive) plan is:</p>
  
  <ul>
  <li>Compute numElementsPerWord = 4 / sizeof(T)</li>
  <li>if(K % numElementsPerWord == 0), then read have thread K read the next full 32-bit word</li>
  <li>store the 32 bit word in shared memory</li>
  <li>after the shared memory has been populated, (and __syncthreads() called) then each thread K can process work on computing output element K </li>
  </ul>
  
  <p>My concern is that it will not coalesce because (for example, in the case where T=short)</p>
  
  <ul>
  <li>Thread 0 reads word 0 from global memory</li>
  <li>Thread 1 does not read </li>
  <li>Thread 2 reads word 1 from global memory</li>
  <li>Thread 3 does not read </li>
  <li>etc...</li>
  </ul>
  
  <p>In other words, thread K reads word (K/sizeof(T)). This would seem to not coalesce properly.</p>
  
  <p>An alternative approach that I considered was:</p>
  
  <ul>
  <li>Launch with number of threads = (N + 3) / 4, such that each thread will be responsible for loading and processing (4/sizeof(T)) elements (each thread processes one 32-bit word - possibly 1, 2, or 4 elements depending on sizeof(T)). However I am concerned that this approach will not be as fast as possible since each thread must then do twice (if T=short) or even quadruple (if T=unsigned char) the amount of processing.</li>
  </ul>
  
  <p>Can someone please tell me if my assumption about my plan is correct: i.e.: it will not coalesce properly?</p>
  
  <p>Can you please comment on my alternative approach?</p>
  
  <p>Can you recommend a more optimal approach that properly coalesces?</p>