OpenCL kernel for working with 3 Bytes blocks
<p>I have a OpenCL kernel, which changes the bit order within 24 bit wide blocks.</p>

<p>My first attempt to realize this is to create <code>buffersize/3</code> threads, but somehow my OpenCl kernel works slower as the same algorithm on the CPU (2 GHz Intel Core i7). </p>

<pre><code>typedef unsigned char uint8_t;

__kernel void decode(
  __global uint8_t* in,
  __global uint8_t* out,
  const unsigned int count) {
    int i = get_global_id(0)*3;

    if(i&lt;count){
        out[i]=in[i];
        out[i+1]=(in[i+1]&amp;0b00001111)&lt;&lt;4|(in[i+2]&amp;0b11110000)&gt;&gt;4;
        out[i+2]=(in[i+2]&amp;0b00001111)&lt;&lt;4|(in[i+1]&amp;0b11110000)&gt;&gt;4;
    }
}
</code></pre>

<p>This kernel is called this way:</p>

<pre><code>count = buffersize/3; // buffersize is approx. 6 to 8 MB
error = clEnqueueNDRangeKernel(
        commands, 
        koDecode, 
        1, 
        NULL, 
        &amp;count, 
        NULL, 
        0, 
        NULL, 
        NULL);
</code></pre>

<p>Is there a better way to do it?</p>