Optimizing opencl kernel <p>I am trying to optimize this kernel. The CPU version of this kernel is 4 times faster than the GPU version. I would expect that the GPU version would be faster. It might be that we have a lot of memory accesses and that is why we have a low performance. I am using an Intel HD 2500 and OpenCL 1.2. </p> <p>The GPU kernel is: </p> <pre><code>__kernel void mykernel(__global unsigned char *inp1, __global unsigned char *inp2, __global unsigned char *inp3, __global unsigned char *inp4, __global unsigned char *outp1, __global unsigned char *outp2, __global unsigned char *outp3, __global unsigned char *outp4, __global unsigned char *lut, uint size ) { unsigned char x1, x2, x3, x4; unsigned char y1, y2, y3, y4; const int x = get_global_id(0); const int y = get_global_id(1); const int width = get_global_size(0); const uint id = y * width + x; x1 = inp1[id]; x2 = inp2[id]; x3 = inp3[id]; x4 = inp4[id]; y1 = (x1 &amp; 0xff) | (x2&gt;&gt;2 &amp; 0xaa) | (x3&gt;&gt;4 &amp; 0x0d) | (x4&gt;&gt;6 &amp; 0x02); y2 = (x1&lt;&lt;2 &amp; 0xff) | (x2 &amp; 0xaa) | (x3&gt;&gt;2 &amp; 0x0d) | (x4&gt;&gt;4 &amp; 0x02); y3 = (x1&lt;&lt;4 &amp; 0xff) | (x2&lt;&lt;2 &amp; 0xaa) | (x3 &amp; 0x0d) | (x4&gt;&gt;2 &amp; 0x02); y4 = (x1&lt;&lt;6 &amp; 0xff) | (x2&lt;&lt;4 &amp; 0xaa) | (x3&lt;&lt;2 &amp; 0x0d) | (x4 &amp; 0x02); // lookup table y1 = lut[y1]; y2 = lut[y2]; y3 = lut[y3]; y4 = lut[y4]; outp1[id] = (y1 &amp; 0xc0) | ((y2 &amp; 0xc0) &gt;&gt; 2) | ((y3 &amp; 0xc0) &gt;&gt; 4) | ((y4 &amp; 0xc0) &gt;&gt; 6); outp2[id] = ((y1 &amp; 0x30) &lt;&lt; 2) | (y2 &amp; 0x30) | ((y3 &amp; 0x30) &gt;&gt; 2) | ((y4 &amp; 0x30) &gt;&gt; 4); outp3[id] = ((y1 &amp; 0x0c) &lt;&lt; 4) | ((y2 &amp; 0x0c) &lt;&lt; 2) | (y3 &amp; 0x0c) | ((y4 &amp; 0x0c) &gt;&gt; 2); outp4[id] = ((y1 &amp; 0x03) &lt;&lt; 6) | ((y2 &amp; 0x03) &lt;&lt; 4) | ((y3 &amp; 0x03) &lt;&lt; 2) | (y4 &amp; 0x03); } </code></pre> <p>I use :</p> <pre><code> size_t localWorkSize[1], globalWorkSize[1]; localWorkSize[0] = 1; globalWorkSize[0] = X*Y; // X,Y define a data space of 15 - 20 MB </code></pre> <p>LocalWorkSize can vary between 1 - 256. </p> <pre><code>for LocalWorkSize = 1 I have CPU = 0.067Sec GPU = 0.20Sec for LocalWorkSize = 256 I have CPU = 0.067Sec GPU = 0.34Sec </code></pre> <p>Which is really weird. Can you give me some ideas why I get these strange numbers? and do you have any tips on how I can optimize this kernel?</p> <p>My main looks like this: </p> <pre><code>int main(int argc, char** argv) { int err,err1,j,i; // error code returned from api calls and other clock_t start, end; // measuring performance variables cl_device_id device_id; // compute device id cl_context context; // compute context cl_command_queue commands; // compute command queue cl_program program_ms_naive; // compute program cl_kernel kernel_ms_naive; // compute kernel // ... dynamically allocate arrays // ... initialize arrays cl_uint dev_cnt = 0; clGetPlatformIDs(0, 0, &amp;dev_cnt); cl_platform_id platform_ids[100]; clGetPlatformIDs(dev_cnt, platform_ids, NULL); // Connect to a compute device err = clGetDeviceIDs(platform_ids[0], CL_DEVICE_TYPE_GPU, 1, &amp;device_id, NULL); // Create a compute context context = clCreateContext(0, 1, &amp;device_id, NULL, NULL, &amp;err); // Create a command queue commands = clCreateCommandQueue(context, device_id, 0, &amp;err); // Create the compute programs from the source file program_ms_naive = clCreateProgramWithSource(context, 1, (const char **) &amp;kernelSource_ms, NULL, &amp;err); // Build the programs executable err = clBuildProgram(program_ms_naive, 0, NULL, NULL, NULL, NULL); // Create the compute kernel in the program we wish to run kernel_ms_naive = clCreateKernel(program_ms_naive, "ms_naive", &amp;err); d_A1 = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, mem_size_cpy/4, h_A1, &amp;err); d_A2 = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, mem_size_cpy/4, h_A2, &amp;err); d_A3 = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, mem_size_cpy/4, h_A3, &amp;err); d_A4 = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, mem_size_cpy/4, h_A4, &amp;err); d_lut = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, 256, h_ltable, &amp;err); d_B1 = clCreateBuffer(context, CL_MEM_WRITE_ONLY, mem_size_cpy/4, NULL, &amp;err); d_B2 = clCreateBuffer(context, CL_MEM_WRITE_ONLY, mem_size_cpy/4, NULL, &amp;err); d_B3 = clCreateBuffer(context, CL_MEM_WRITE_ONLY, mem_size_cpy/4, NULL, &amp;err); d_B4 = clCreateBuffer(context, CL_MEM_WRITE_ONLY, mem_size_cpy/4, NULL, &amp;err); int size = YCOLUMNS*XROWS/4; int size_b = size * 4; err = clSetKernelArg(kernel_ms_naive, 0, sizeof(cl_mem), (void *)&amp;(d_A1)); err |= clSetKernelArg(kernel_ms_naive, 1, sizeof(cl_mem), (void *)&amp;(d_A2)); err |= clSetKernelArg(kernel_ms_naive, 2, sizeof(cl_mem), (void *)&amp;(d_A3)); err |= clSetKernelArg(kernel_ms_naive, 3, sizeof(cl_mem), (void *)&amp;(d_A4)); err |= clSetKernelArg(kernel_ms_naive, 4, sizeof(cl_mem), (void *)&amp;d_B1); err |= clSetKernelArg(kernel_ms_naive, 5, sizeof(cl_mem), (void *)&amp;(d_B2)); err |= clSetKernelArg(kernel_ms_naive, 6, sizeof(cl_mem), (void *)&amp;(d_B3)); err |= clSetKernelArg(kernel_ms_naive, 7, sizeof(cl_mem), (void *)&amp;(d_B4)); err |= clSetKernelArg(kernel_ms_naive, 8, sizeof(cl_mem), (void *)&amp;d_lut); //__global err |= clSetKernelArg(kernel_ms_naive, 9, sizeof(cl_uint), (void *)&amp;size_b); size_t localWorkSize[1], globalWorkSize[1]; localWorkSize[0] = 256; globalWorkSize[0] = XROWS*YCOLUMNS; start = clock(); for (i=0;i&lt; EXECUTION_TIMES;i++) { err1 = clEnqueueNDRangeKernel(commands, kernel_ms_naive, 1, NULL, globalWorkSize, localWorkSize, 0, NULL, NULL); err = clFinish(commands); } end = clock(); return 0; } </code></pre>