title: OpenCL Kernel Executes Slower Than Single Thread
body: <p>All, I wrote a very simple OpenCL kernel which transforms an RGB image to gray scale using simple averaging.</p>
  
  <p>Some background:</p>
  
  <ol>
  <li>The image is stored in mapped memory, as a 24 bit, non padded memory block</li>
  <li>The output array is stored in pinned memory (mapped with <code>clEnqueueMapBuffer</code>) and is 8 bpp</li>
  <li>There are two buffers allocated on the device (<code>clCreateBuffer</code>), one is specifically read (which we <code>clWriteBuffer</code> into before the kernel starts) and the other is specifically write (which we <code>clReadBuffer</code> after the kernel finishes)</li>
  </ol>
  
  <p>I am running this on a 1280x960 image. A serial version of the algorithm averages 60ms, the OpenCL kernel averages 200ms!!! I'm doing something wrong but I have no idea how to proceed, what to optimize. (Timing my reads/writes without a kernel call, the algorithm runs in 15ms)</p>
  
  <p>I am attaching the kernel setup (sizes and arguments) as well as the kernel</p>
  
  <hr>
  
  <p><strong>EDIT</strong>: So I wrote an even dumber kernel, that does no global memory accesses inside it, and it was only 150ms... This is still ridiculously slow. I thought maybe I'm messing up with global memory reads, they have to be 4 byte aligned or something? Nope...</p>
  
  <p><strong>Edit 2:</strong> Removing the all parameters from my kernel gave me significant speed up... I'm confused I thought that since I'm <code>clEnqueueWriteBuffer</code> the kernel should be doing no memory transfer from host->device and device->host....</p>
  
  <p><strong>Edit 3:</strong> Figured it out, but I still don't understand why. If anyone could explain it I would be glad to award correct answer to them. The problem was passing the custom structs by value. It looks like I'll need to allocate a global memory location for them and pass their <code>cl_mem</code>s</p>
  
  <hr>
  
  <p>Kernel Call:</p>
  
  <pre><code>//Copy input to device
  result = clEnqueueWriteBuffer(handles-&gt;queue, d_input_data, CL_TRUE, 0, h_input.widthStep*h_input.height, (void *)input-&gt;imageData, 0, 0, 0);
  if(check_result(result, "opencl_rgb_to_gray", "Failed to write to input buffer on device!")) return 0;
  
  //Set kernel arguments
  result = clSetKernelArg(handles-&gt;current_kernel, 0, sizeof(OpenCLImage), (void *)&amp;h_input);
  if(check_result(result, "opencl_rgb_to_gray", "Failed to set input struct.")) return 0;
  result = clSetKernelArg(handles-&gt;current_kernel, 1, sizeof(cl_mem), (void *)&amp;d_input_data);
  if(check_result(result, "opencl_rgb_to_gray", "Failed to set input data.")) return 0;
  result = clSetKernelArg(handles-&gt;current_kernel, 2, sizeof(OpenCLImage), (void *)&amp;h_output);
  if(check_result(result, "opencl_rgb_to_gray", "Failed to set output struct.")) return 0;
  result = clSetKernelArg(handles-&gt;current_kernel, 3, sizeof(cl_mem), (void *)&amp;d_output_data);
  if(check_result(result, "opencl_rgb_to_gray", "Failed to set output data.")) return 0;
  
  //Determine run parameters
  global_work_size[0] = input-&gt;width;//(unsigned int)((input-&gt;width / (float)local_work_size[0]) + 0.5);
  global_work_size[1] = input-&gt;height;//(unsigned int)((input-&gt;height/ (float)local_work_size[1]) + 0.5);
  
  printf("Global Work Group Size: %d %d\n", global_work_size[0], global_work_size[1]);
  
  //Call kernel
  result = clEnqueueNDRangeKernel(handles-&gt;queue, handles-&gt;current_kernel, 2, 0, global_work_size, local_work_size, 0, 0, 0);
  if(check_result(result, "opencl_rgb_to_gray", "Failed to run kernel!")) return 0;
  
  result = clFinish(handles-&gt;queue);
  if(check_result(result, "opencl_rgb_to_gray", "Failed to finish!")) return 0;
  
  //Copy output
  result = clEnqueueReadBuffer(handles-&gt;queue, d_output_data, CL_TRUE, 0, h_output.widthStep*h_output.height, (void *)output-&gt;imageData, 0, 0, 0);
  if(check_result(result, "opencl_rgb_to_gray", "Failed to write to output buffer on device!")) return 0;
  </code></pre>
  
  <p>Kernel:</p>
  
  <pre><code>typedef struct OpenCLImage_t
  {
      int width;
      int widthStep;
      int height;
      int channels;
  } OpenCLImage;
  
  __kernel void opencl_rgb_kernel(OpenCLImage input, __global unsigned char*  input_data, OpenCLImage output, __global unsigned char * output_data)
  {
      int pixel_x = get_global_id(0);
      int pixel_y = get_global_id(1);
      unsigned char * cur_in_pixel, *cur_out_pixel;
      float avg = 0;
  
      cur_in_pixel = (unsigned char *)(input_data + pixel_y*input.widthStep + pixel_x * input.channels);
      cur_out_pixel = (unsigned char *)(output_data + pixel_y*output.widthStep + pixel_x * output.channels);
  
      avg += cur_in_pixel[0];
      avg += cur_in_pixel[1];
      avg+= cur_in_pixel[2];
      avg /=3.0f;
  
      if(avg &gt; 255.0)
          avg = 255.0;
      else if(avg &lt; 0)
          avg = 0;
  
      *cur_out_pixel = avg;
  }
  </code></pre>