title: About compact operation in cuddpp
body: <p>The following kernel function is the compact operation in the cudpp, a cuda library (http://gpgpu.org/developer/cudpp).</p>
  
  <p>My question is why the developer repeats the the writing part 8 times? And why it can improve the performance?</p>
  
  <p>And why one thread process 8 elements, why not each thread process one element?</p>
  
  <pre><code>template &lt;class T, bool isBackward&gt;
  __global__ void compactData(T                  *d_out, 
                          size_t             *d_numValidElements,
                          const unsigned int *d_indices, // Exclusive Sum-Scan Result
                          const unsigned int *d_isValid,
                          const T            *d_in,
                          unsigned int       numElements)
  {
    if (threadIdx.x == 0)
    {
          if (isBackward)
              d_numValidElements[0] = d_isValid[0] + d_indices[0];
      else
          d_numValidElements[0] = d_isValid[numElements-1] + d_indices[numElements-1];
     }
  
     // The index of the first element (in a set of eight) that this
     // thread is going to set the flag for. We left shift
     // blockDim.x by 3 since (multiply by 8) since each block of 
     // threads processes eight times the number of threads in that
     // block
     unsigned int iGlobal = blockIdx.x * (blockDim.x &lt;&lt; 3) + threadIdx.x;
  
     // Repeat the following 8 (SCAN_ELTS_PER_THREAD) times
     // 1. Check if data in input array d_in is null
     // 2. If yes do nothing
     // 3. If not write data to output data array d_out in
     //    the position specified by d_isValid
     if (iGlobal &lt; numElements &amp;&amp; d_isValid[iGlobal] &gt; 0) {
         d_out[d_indices[iGlobal]] = d_in[iGlobal];
     }
     iGlobal += blockDim.x;  
     if (iGlobal &lt; numElements &amp;&amp; d_isValid[iGlobal] &gt; 0) {
         d_out[d_indices[iGlobal]] = d_in[iGlobal];       
     }
     iGlobal += blockDim.x;
     if (iGlobal &lt; numElements &amp;&amp; d_isValid[iGlobal] &gt; 0) {
         d_out[d_indices[iGlobal]] = d_in[iGlobal];
     }
     iGlobal += blockDim.x;
     if (iGlobal &lt; numElements &amp;&amp; d_isValid[iGlobal] &gt; 0) {
         d_out[d_indices[iGlobal]] = d_in[iGlobal];
     }
     iGlobal += blockDim.x;
     if (iGlobal &lt; numElements &amp;&amp; d_isValid[iGlobal] &gt; 0) {
         d_out[d_indices[iGlobal]] = d_in[iGlobal];
     }
     iGlobal += blockDim.x;
     if (iGlobal &lt; numElements &amp;&amp; d_isValid[iGlobal] &gt; 0) {
         d_out[d_indices[iGlobal]] = d_in[iGlobal];
     }
     iGlobal += blockDim.x;
     if (iGlobal &lt; numElements &amp;&amp; d_isValid[iGlobal] &gt; 0) {
         d_out[d_indices[iGlobal]] = d_in[iGlobal];
     }
     iGlobal += blockDim.x;
     if (iGlobal &lt; numElements &amp;&amp; d_isValid[iGlobal] &gt; 0) {
         d_out[d_indices[iGlobal]] = d_in[iGlobal];
     }
  }
  </code></pre>