title: Optimize Cuda Kernel time execution
body: <p>I'm a learning Cuda student, and I would like to optimize the execution time of my kernel function. As a result, I realized a short program computing the difference between two pictures. So I compared the execution time between a classic CPU execution in C, and a GPU execution in Cuda C.</p>
  
  <p>Here you can find the code I'm talking about:</p>
  
  <pre><code>int *imgresult_data = (int *) malloc(width*height*sizeof(int));
  int size = width*height;
  
  switch(computing_type)
  {
  
      case GPU:
  
      HANDLE_ERROR(cudaMalloc((void**)&amp;dev_data1, size*sizeof(unsigned char)));
      HANDLE_ERROR(cudaMalloc((void**)&amp;dev_data2, size*sizeof(unsigned char)));
      HANDLE_ERROR(cudaMalloc((void**)&amp;dev_data_res, size*sizeof(int)));
  
      HANDLE_ERROR(cudaMemcpy(dev_data1, img1_data, size*sizeof(unsigned char), cudaMemcpyHostToDevice)); 
      HANDLE_ERROR(cudaMemcpy(dev_data2, img2_data, size*sizeof(unsigned char), cudaMemcpyHostToDevice));
      HANDLE_ERROR(cudaMemcpy(dev_data_res, imgresult_data, size*sizeof(int), cudaMemcpyHostToDevice));
  
      float time;
      cudaEvent_t start, stop;
  
      HANDLE_ERROR( cudaEventCreate(&amp;start) );
      HANDLE_ERROR( cudaEventCreate(&amp;stop) );
      HANDLE_ERROR( cudaEventRecord(start, 0) );
  
      for(int m = 0; m &lt; nb_loops ; m++)
      {
          diff&lt;&lt;&lt;height, width&gt;&gt;&gt;(dev_data1, dev_data2, dev_data_res);
      }
  
      HANDLE_ERROR( cudaEventRecord(stop, 0) );
      HANDLE_ERROR( cudaEventSynchronize(stop) );
      HANDLE_ERROR( cudaEventElapsedTime(&amp;time, start, stop) );
  
      HANDLE_ERROR(cudaMemcpy(imgresult_data, dev_data_res, size*sizeof(int), cudaMemcpyDeviceToHost));
  
      printf("Time to generate:  %4.4f ms \n", time/nb_loops);
  
      break;
  
      case CPU:
  
      clock_t begin = clock(), diff;
  
      for (int z=0; z&lt;nb_loops; z++)
      {
          // Apply the difference between 2 images
          for (int i = 0; i &lt; height; i++)
          {
              tmp = i*imgresult_pitch;
              for (int j = 0; j &lt; width; j++)
              {
                  imgresult_data[j + tmp] = (int) img2_data[j + tmp] - (int) img1_data[j + tmp];
              }
          }
      }
      diff = clock() - begin;
  
      float msec = diff*1000/CLOCKS_PER_SEC;
      msec = msec/nb_loops;
      printf("Time taken %4.4f milliseconds", msec);
  
      break;
  }
  </code></pre>
  
  <p>And here is my kernel function:</p>
  
  <pre><code>__global__ void diff(unsigned char *data1 ,unsigned char *data2, int *data_res)
  {
      int row = blockIdx.x;
      int col = threadIdx.x;
      int v = col + row*blockDim.x;
  
      if (row &lt; MAX_H &amp;&amp; col &lt; MAX_W)
      {
          data_res[v] = (int) data2[v] - (int) data1[v];
      }
  }
  </code></pre>
  
  <p>I obtained these execution time for each one</p>
  
  <ul>
  <li>CPU: 1,3210ms</li>
  <li>GPU: 0,3229ms</li>
  </ul>
  
  <p>I wonder why GPU result is not as lower as it should be. I am a beginner in Cuda so please be comprehensive if there are some classic errors.</p>
  
  <p>EDIT1:
  Thank you for your feedback. I tried to delete the 'if' condition from the kernel but it didn't change deeply my program execution time.</p>
  
  <p>However, after having install Cuda profiler, it told me that my threads weren't running concurrently. I don't understand why I have this kind of message, but it seems true because I only have a 5 or 6 times faster application with GPU than with CPU. This ratio should be greater, because each thread is supposed to process one pixel concurrently to all the other ones. If you have an idea of what I am doing wrong, it would be hepful...</p>
  
  <p>Flow.</p>