title: How to merge row processing kernel filterring and column processing kernel filterring into single openCL kernel
body: <p>I have image processing filter implemented using row and column processing in openCL, Where most of the processing gets wasted in launching multiple kernels itself.</p>
  
  <p>So Can i merge these two kernel into single kernel which do the same functionality and performs better in Intel HD4600 graphics card. Details of the code is given below:-</p>
  
  <p>Assumptions:<br> 
  1. Both horizontal and vertical padding is done by host (c programming)<br>
  2. N (filter length is 8, width and height are 1024 x 1024,filter coefficients are generated using generic filters <br>
  3. First row and then col Kernel is being launched using  below API<br> 
  ret |= clEnqueueNDRangeKernel(command_queue, kernel, 2, NULL, global_ws(1024x1024), NULL, 0, NULL,NULL);<br></p>
  
  <p>//Code:<br></p>
  
  <pre><code>__kernel void filter_rows(__global float *ip_img,__global float *op_img,
                            int width, int height,int pitch,int N,__constant float *W)
  {
      __private int i=get_global_id(0); 
      __private int j=get_global_id(1); 
      __private int k;
      __private float a;
      __private int image_offset = N*pitch +N;
      __private int curr_pix = j*pitch + i +image_offset;
      // apply filter
      for(k=-N, a=0.0f; k&lt;=N; k++)
      {
      a += ip_img[curr_pix+k] * W[k+N];
      }
      op_img[curr_pix] = a;   
  }
  __kernel void filter_col(__global float *ip_img,__global float *op_img,int width,
                           int height,int pitch,int N,__constant float *W)
  {
      __private int i=get_global_id(0);
      __private int j=get_global_id(1);
      __private int k;
      __private float a;
      __private int image_offset = N*pitch +N;
      __private int curr_pix = j*pitch + i +image_offset;
  
      // apply filter
      for(k=-N, a=0.0f; k&lt;=N; k++)
      {
        a += ip_img[k*pitch +curr_pix] * W[k+N];                  
      }
      op_img[curr_pix] = a;
  }
  void padd_hor(float *ip_img,pad_leng)
  {
      //...using simple C programming
  }
  void padd_ver(float *ip_img,pad_leng)
  {
      //...using simple c programming
  }
  void generic_filter(_global float *in_image,__global float *out_image,
  __global float *temp_image,int width, int height,int pitch,int N,
  __constant float *Wr,__constant float *Wc)
  {
      padd_hor(in_image,filter_length)
      filter_rows(in_image,temp_image,width,height,pitch,filter_length,filter_coeff_hor);
      pad_ver(temp_image,filter_length)
      filter_col(temp_image,out_image,width,height,pitch,filter_length,filter_coeff_ver);
  }
  __kernel generic_filter(_global float *in_image,__global float *out_image,__global      float*temp_image,
  int width, int height,int pitch,int N,__constant float *Wr,__constant float *Wc)
  {
      // ... here i need your suggetion to implement the kernel which do same as generic_filter
  }
  </code></pre>
  
  <p>Your help will be appreciated to optimize this filter and  get the best possible result.<br>
  Also please let me know how much max gain can i  get with respect to C code runs on Intel CPU.<br></p>
  
  <p>Thanks and Regards <br>
  Vijayky88</p>