title: How can I optimize my OpenCL kernel for the GPU?
body: <p>This question follows from <a href="https://stackoverflow.com/questions/45927692/how-can-i-make-idct-run-faster-on-my-gpu?noredirect=1#comment78844616_45927692">this question</a>. <strong>Since then, I have been working on my IDCT kernel to make it more efficient.</strong></p>
  
  <p>This version of the kernel produces the correct output but is slow:</p>
  
  <pre><code>void idct_1D(__local int *Y);
  
   __kernel void IDCT(__global int* input, __global uchar* output) 
  {
   unsigned int kid= get_global_id(0);
  
   __local int Y[64]; 
   int k,l;
   __local int Yc[8];
  
   for (k = 0; k &lt; 8; k++)
   {
    for (l = 0; l &lt; 8; l++)
    {
     Y(k,l) = SCALE(input[(k &lt;&lt; 3) + l], S_BITS);     
    }         
   idct_1D(&amp;Y(k,0));
   }
  
  for (l = 0; l &lt; 8; l++)
  {
  for (k = 0; k &lt; 8; k++)
  {Yc[k] = Y(k, l);}
  
  idct_1D(Yc);
  
  for (k = 0; k &lt; 8; k++)
  {
  
  int r = 128 + DESCALE(Yc[k], S_BITS + 3);
  r = r &gt; 0 ? (r &lt; 255 ? r : 255) : 0;
  X(k, l) = r;
  }
  
  }
  
  }
  </code></pre>
  
  <p>I tried to make this kernel more parallel by modifying its structure in this way:</p>
  
  <pre><code>__kernel void IDCT(__global int* input, __global uchar* output) 
  {
  unsigned int kid= get_global_id(0);
  
      __local int Y[64]; 
      int k= get_global_id(0);
      int l;
      int lid= get_global_id(1);
      __local int Yc[8];
  
     if (k &lt; 8)
      {
          for (l = 0; l &lt; 8; l++) 
  
       {
       Y(k, l) = SCALE(input[(k &lt;&lt; 3) + l], S_BITS);
       }
          idct_1D(&amp;Y(k, 0));
      }
  
      if (lid &lt; 8)
      {
  
          for (k = 0; k &lt; 8; k++)
      {
              Yc[k] = Y(k, lid);
      }
  
          idct_1D(Yc);
  
          for (k = 0; k &lt; 8; k++)
          {
              int r = 128 + DESCALE(Yc[k], S_BITS + 3);
              r = r &gt; 0 ? (r &lt; 255 ? r : 255) : 0;
              X(k, lid) = r;
          }
  
      }
  }
  </code></pre>
  
  <p>The kernel above is giving me the correct output but I am seeing no changes in the processing speed of my code. </p>
  
  <p>My <code>main.c</code> where I am calling the code from looks like this:</p>
  
  <pre><code>  for (index_X = 0; index_X &lt; nb_MCU_X; index_X++) {
  
      for (index_Y = 0; index_Y &lt; nb_MCU_Y; index_Y++) {
  
      for (index = 0; index &lt; SOS_section.n; index++) {
  
       uint32_t component_index = component_order[index];
  
       int nb_MCU = ((SOF_component[component_index].HV&gt;&gt; 4) &amp; 0xf) * (SOF_component[component_index].HV &amp; 0x0f);
  
       for (chroma_ss = 0; chroma_ss &lt; nb_MCU; chroma_ss++) {
  
      cl_mem DCT_Input = clCreateBuffer(context, CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR, 64 * sizeof(cl_int), unZZ_MCU, &amp;ret);
  
      //Output buffer
      cl_mem  DCT_Output = clCreateBuffer(context, CL_MEM_READ_WRITE| CL_MEM_COPY_HOST_PTR, (MCU_sx * MCU_sy * max_ss_h * max_ss_v) + 4, YCbCr_MCU_ds[component_index] + (64 * chroma_ss), &amp;ret);
      chk(ret, "clCreateBuffer");
  
      ret = clSetKernelArg(cos_kernel, 0, sizeof(cl_mem), (void *)&amp;DCT_Input);
      ret |= clSetKernelArg(cos_kernel, 1, sizeof(cl_mem), (void *)&amp;DCT_Output);
  
  
  const size_t globalForInverseDCT[2]= {8, 8};
  
  
  
   ret = clEnqueueNDRangeKernel(command_queue, cos_kernel, 2, NULL, &amp;globalForInverseDCT, NULL, 0, NULL, NULL);
  
  //Timing-End..
  
   ret = clEnqueueReadBuffer(command_queue, DCT_Output, CL_TRUE, 0, (MCU_sx * MCU_sy * max_ss_h * max_ss_v) + 4, YCbCr_MCU_ds[component_index] + (64 * chroma_ss), 0, NULL, NULL);
  
   }
  
   //other function
  
   }
  
   //code continues...
  </code></pre>
  
  <p><strong>How can I further optimize this kernel?</strong></p>
  
  <p>EDIT:</p>
  
  <p>Ideally, I wish to break <code>IDCT</code> into even smaller kernels. Thanks to the help I received in <a href="https://stackoverflow.com/questions/45901912/how-can-i-split-this-function-into-smaller-functions-without-sacrificing-its-fun/45903658#45903658">this question</a>, the regular function was broken down into smaller sub-functions like this:</p>
  
  <pre><code>void IDCTforX(int32_t *input, uint8_t *output) {
  
      int32_t Y[64];
      int32_t k, l;
      int32_t Yc[8];
  
      for (k = 0; k &lt; 8; k++) {
          for (l = 0; l &lt; 8; l++)
          {
              Y(k, l) = SCALE(input[(k &lt;&lt; 3) + l], S_BITS);
          }
      }
  }
  
  void IDCTfor1dim(int32_t *input, uint8_t *output)
  {
  int32_t Y[64];
      int32_t k, l;
      int32_t Yc[8];
  
      for (k= 0; k &lt; 8; k++)
      {
          idct_1d(&amp;Y(k, 0));
      }
  }
  </code></pre>
  
  <p>In the Y-direction:</p>
  
  <pre><code>void IDCTforY(int32_t *input, uint8_t *output) 
  {
    int32_t Y[64];
    int32_t k, l;
    int32_t Yc[8][8];
  
    for (l = 0; l &lt; 8; l++) 
    {
        for (k = 0; k &lt; 8; k++)
            Yc[l][k] = Y(k, l);
        idct_1d(Yc[l]);
     }
  
  void IDCT_DescaleY(int32_t *input, uint8_t *output) 
  {
    int32_t Y[64];
    int32_t k, l;
    int32_t Yc[8][8];
  
  for (l = 0; l &lt; 8; l++) 
     {
         for (k = 0; k &lt; 8; k++) 
         {   
             int32_t r = 128 + DESCALE(Yc[l][k], S_BITS + 3);
             r = r &gt; 0 ? (r &lt; 255 ? r : 255) : 0;
             X(k, l) = r;
         }
     }
  }
  </code></pre>
  
  <p>In <code>main.c</code>, I get the correct output when I call the functions (in place of the original function) in this sequence:</p>
  
  <pre><code>IDCTforX(unZZ_MCU, YCbCr_MCU_ds[component_index] + (64 * chroma_ss));
  IDCTfor1dim(unZZ_MCU, YCbCr_MCU_ds[component_index] + (64 * chroma_ss));
  IDCTforY(unZZ_MCU, YCbCr_MCU_ds[component_index] + (64 * chroma_ss));
  IDCT_DescaleY(unZZ_MCU, YCbCr_MCU_ds[component_index] + (64 * chroma_ss));
  </code></pre>
  
  <p>I translated these functions into kernels but I am only getting a gray screen. My new kernels look like this:</p>
  
  <pre><code>/*---------------IDCTForX----------------------------*/
  
  __kernel void IDCTforX(__global int *input, __global uchar *output) {
  
      int Y[64];
      unsigned int k= get_global_id(0);
      unsigned int l= get_global_id(1);
  
          if ((k &lt; 8) &amp;&amp; (l &lt; 8))
              {
                  Y(k, l) = SCALE(input[(k &lt;&lt; 3) + l], S_BITS);
  
              }
  
  }
  
  /*---------------IDCTfor1Dim-------------------------*/
  
  __kernel void IDCTfor1dim(__global int *input, __global uchar *output){
  
  int Y[64];
  unsigned int k= get_global_id(0);
  
      for (k= 0; k &lt; 8; k++)
      {
          idct_1D(&amp;Y(k, 0));
      }
  
  }
  
  /*---------------IDCTYSplit----------------------------*/
  
  __kernel void IDCTYSplit(__global int *input, __global uchar *output) {
  
  int Y[64];
        int k= get_global_id(0);
        int l= get_global_id(1);
        int Yc[8][8];
  
        if ((k &lt; 8) &amp;&amp; (l &lt; 8))
        {
                Yc[l][k] = Y(k, l);
  
        }
  }
  
  /*---------------IDCTY_Inverse----------------------------*/
  
  __kernel void IDCTY_Inverse(__global int *input, __global uchar *output) {
  
  int Y[64];
        int l= get_global_id(0);
        int Yc[8][8];
  
   if (l &lt; 8)
            {
            idct_1D(Yc[l]);
         }
  }
  
  /*---------------IDCTY_Descale----------------------------*/
  
  __kernel void IDCTY_Descale(__global int *input, __global uchar *output) {
  
  int Y[64];
  
      int l= get_global_id(0);
  
      int k= get_global_id(1);
  
        int Yc[8][8];
  
         if ((l &lt; 8) &amp;&amp; (k &lt; 8))
             {
                 int r = 128 + DESCALE(Yc[l][k], S_BITS + 3);
                 r = r &gt; 0 ? (r &lt; 255 ? r : 255) : 0;
                 X(k, l) = r;
             }
  
  }
  </code></pre>