How to pack bits (efficiently) in CUDA?
<p>I have an array of bytes where each byte is either 0 or 1. Now I want to pack these values into bits, so that 8 original bytes occupy 1 target byte, with original byte 0 going into bit 0, byte 1 into bit 1, etc.
So far I have the following in the kernel:</p>

<pre><code>const uint16_t tid = threadIdx.x;
__shared__ uint8_t packing[cBlockSize];

// ... Computation of the original bytes in packing[tid]
__syncthreads();

if ((tid &amp; 4) == 0)
{
    packing[tid] |= packing[tid | 4] &lt;&lt; 4;
}
if ((tid &amp; 6) == 0)
{
    packing[tid] |= packing[tid | 2] &lt;&lt; 2;
}
if ((tid &amp; 7) == 0)
{
    pOutput[(tid + blockDim.x*blockIdx.x)&gt;&gt;3] = packing[tid] | (packing[tid | 1] &lt;&lt; 1);
}
</code></pre>

<p>Is this correct and efficient?</p>