title: How can I make IDCT run faster on my GPU?
body: <p>I am trying to optimize IDCT from <a href="https://github.com/claf/MJPEG/tree/master/tima_seq_version/src" rel="nofollow noreferrer">this</a> code for the GPU. The GPU I have on my system in <strong>NVIDIA Tesla k20c</strong>.</p>
  
  <p>The <code>IDCT</code> function as written in the original code looks like this:</p>
  
  <pre><code>void IDCT(int32_t *input, uint8_t *output) {
      int32_t Y[64];
      int32_t k, l;
  
      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++) {
          int32_t Yc[8];
  
          for (k = 0; k &lt; 8; k++) Yc[k] = Y(k, l);
  
          idct_1d(Yc);
  
          for (k = 0; k &lt; 8; k++) {
              int32_t 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>More details about the <code>.c</code> file can be found in <a href="https://github.com/claf/MJPEG/blob/master/tima_seq_version/src/idct.c" rel="nofollow noreferrer">this</a> link.</p>
  
  <p>This function is called from <code>main.c</code> in the following way:</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++) {
  
   //unpack block function
   //iqzz function
  
  IDCT(unZZ_MCU, YCbCr_MCU_ds[component_index] + (64 * chroma_ss));
  
   }
  //other function                    
   }
  //Code continues...
  </code></pre>
  
  <p>Details of <code>main.c</code> can be found in <a href="https://github.com/claf/MJPEG/blob/master/tima_seq_version/src/main.c" rel="nofollow noreferrer">this</a> link.</p>
  
  <p>I made a kernel out of the IDCT function in the following way:</p>
  
  <pre><code>__kernel void IDCT(__global int* input, __global uchar* output) 
  {
   unsigned int kid= get_global_id(0);
  
   int Y[64]; 
   int k,l;
   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 am calling this kernel from <code>main.c</code> in this way:</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);
  
  //Timing-Start..
  
    start_time = wtime();
  
    size_t globalForInverseDCT= 1024;
  
    size_t localForInverseDCT= 256;
  
   ret = clEnqueueNDRangeKernel(command_queue, cos_kernel, 1, NULL, &amp;globalForInverseDCT, &amp;localForInverseDCT, 0, NULL, NULL);
  
  //Timing-End..
  
   run_time = wtime() - start_time;
  
   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>I am using this function for timing <code>IDCT</code>:</p>
  
  <pre><code>double wtime()
  {
  
     /* Use a generic timer */
     static int sec = -1;
     struct timeval tv;
     gettimeofday(&amp;tv, NULL);
     if (sec &lt; 0) sec = tv.tv_sec;
     return (tv.tv_sec - sec) + 1.0e-6*tv.tv_usec;
  
  }
  </code></pre>
  
  <p>The execution time for IDCT on the <strong>CPU</strong> with my timing method is <strong>5 microseconds</strong>. The execution time on the <strong>GPU</strong> is <strong>20 microseconds</strong>. (I used OpenCL's built-in profiling command as well and it gave me an execution time of 31 microseconds.)</p>
  
  <p><strong>How can I modify my code so that it runs faster on the GPU?</strong></p>
  
  <p><strong>MY WORK:</strong></p>
  
  <p>I tried to break my kernel down into smaller kernels in <a href="https://www.jdoodle.com/embed/v0/c/gcc-5.3.0/5Ss" rel="nofollow noreferrer">this</a> way but I am getting static as my screen output.</p>
  
  <p>I profiled the code again and I got the following output:</p>
  
  <p><a href="https://i.stack.imgur.com/06gnj.jpg" rel="nofollow noreferrer"><img src="https://i.stack.imgur.com/06gnj.jpg" alt="enter image description here"></a></p>
  
  <p>The execution time for my IDCT kernel according to Nvidia's Visual Profiler is <strong>17.6 microseconds</strong> (approx.).</p>
  
  <p><strong>EDIT:</strong></p>
  
  <p>I profiled my code again with the command line profiler and the execution time for idct is now <strong>4.9 microseconds</strong>.</p>
  
  <p><strong>PLEASE SUGGEST HOW I CAN FURTHER OPTIMIZE MY KERNEL.</strong></p>