title: CPU overtakes the GPU
body: <p>I have some doubts about understanding the CUDA thread processing in the SM. The following propositions are inferred from what I have been reading:
  My GPU is: GTX650Ti. </p>
  
  <ol>
  <li>Thread count in a block must be ALWAYS a multiple of the Warp size. So, each SM can process blocks of 32 threads (warpSize).</li>
  <li>The maximum thread count my SM can compute at same time is 2048 (maxThreadsPerMultiProcessor). </li>
  <li>Due to 2048 threads can be computed at same time in each SM and the warpSize is 32, then 64 blocks can be computed at same time. </li>
  <li>Due to my GPU has 4 SMs, there can be 64X4=256 blocks of threads executed at same time. </li>
  <li>Therefore, the kernel launch may have the following launch parameters: &lt;&lt;&lt;256, 32>>> and each kernel launch will invoke 8192 threads.</li>
  </ol>
  
  <p>Is that right?</p>
  
  <p>So if I have a vector of 10M elements to process in my kernel it means that I have to segment it in 1221 jobs (kernel launches) of 8192 elements each?</p>
  
  <p>This quest arised because I am comparing the time performance between a sequential program and my CUDA program. But all I can see is that the CPU overtakes the GPU. I also tried with the maximum launch parameters such as &lt;&lt;&lt;65535, 1024>>>. The results are very similar. </p>
  
  <p>So, what am I doing or configuring wrong? </p>
  
  <p>This is the code I'm using:</p>
  
  <pre><code>#include "cuda_runtime.h"
  #include "device_launch_parameters.h"
  #include &lt;math.h&gt;
  #include &lt;time.h&gt;
  #include "C:\cdev.h"
  #include &lt;thrust/device_vector.h&gt;
  
  using namespace thrust;
  using namespace std;
  
  #define N (1024 * 16384)
  
  cdev devices;
  
  __global__ void eucliDist(double *c, double *a, double *b)
  {
      int i = blockDim.x * blockIdx.x + threadIdx.x;
      if (i &lt; N)
          c[i] = sqrt(pow(a[i], 2) + pow(b[i], 2));
  }
  
  int main()
  {
      clock_t start, end;
      double elapsed;
      static double A[N];
      static double B[N];
      for (int i = 0; i &lt; N; i++)
      {
          A[i] = double(i);
          B[i] = double(i);
      }   
      static double C[N];
  
      // Sequential execution of F(x,y) = sqrt((x^2 + y^2))
      start = clock();
      for (int i = 0; i &lt; N; i++)
          C[i] = sqrt(pow(A[i], 2) + pow(B[i], 2));
      end = clock();
      elapsed = double(end - start) / CLOCKS_PER_SEC;
      cout &lt;&lt; "Elapsed time for sequential processing is: " &lt;&lt; elapsed &lt;&lt; " seconds." &lt;&lt; endl;
  
      // CUDA Initialization
      unsigned int threadNum;
      unsigned int blockNum;
      cudaError_t cudaStatus;
      threadNum = devices.ID[0].maxThreadsPerBlock;
      blockNum = ceil(double(N) / double(threadNum));
      // Parallel execution with Thrust of F(x,y) = sqrt((x^2 + y^2))
      vector&lt;double&gt; vectorA(N);
      vector&lt;double&gt; vectorB(N);
      for (int i = 0; i &lt; N; i++)
      {
          vectorA[i] = double(i);
          vectorB[i] = double(i);
      }
      vector&lt;double&gt; vectorC(N);
      start = clock();
      device_vector&lt;double&gt; thrustA(N);
      cudaStatus = cudaGetLastError();
      if (cudaStatus != cudaSuccess)
      {
          cerr &lt;&lt; "Device vector allocation failed: " &lt;&lt; cudaGetErrorString(cudaStatus) &lt;&lt; " (thrustA)" &lt;&lt; endl;
          cin.get();
          return 1;
      }
      device_vector&lt;double&gt; thrustB(N);
      cudaStatus = cudaGetLastError();
      if (cudaStatus != cudaSuccess)
      {
          cerr &lt;&lt; "Device vector allocation failed: " &lt;&lt; cudaGetErrorString(cudaStatus) &lt;&lt; " (thrustB)" &lt;&lt; endl;
          cin.get();
          return 1;
      }
      device_vector&lt;double&gt; thrustC(N);
      cudaStatus = cudaGetLastError();
      if (cudaStatus != cudaSuccess)
      {
          cerr &lt;&lt; "Device vector allocation failed: " &lt;&lt; cudaGetErrorString(cudaStatus) &lt;&lt; " (thrustC)" &lt;&lt; endl;
          cin.get();
          return 1;
      }
      thrustA = vectorA;
      cudaStatus = cudaGetLastError();
      if (cudaStatus != cudaSuccess)
      {
          cerr &lt;&lt; "Host to device copy failed (Thrust): " &lt;&lt; cudaGetErrorString(cudaStatus) &lt;&lt; " (vectorA -&gt; thrustA)" &lt;&lt; endl;
          cin.get();
          return 1;
      }
      thrustB = vectorB;
      cudaStatus = cudaGetLastError();
      if (cudaStatus != cudaSuccess)
      {
          cerr &lt;&lt; "Host to device copy failed (Thrust): " &lt;&lt; cudaGetErrorString(cudaStatus) &lt;&lt; " (vectorB -&gt; thrustB)" &lt;&lt; endl;
          cin.get();
          return 1;
      }
      eucliDist &lt;&lt;&lt;blockNum, threadNum&gt;&gt;&gt;(raw_pointer_cast(thrustC.data()), raw_pointer_cast(thrustA.data()), raw_pointer_cast(thrustB.data()));
      cudaStatus = cudaGetLastError();
      if (cudaStatus != cudaSuccess)
      {
          cerr &lt;&lt; "Kernel launch failed (Thrust): " &lt;&lt; cudaGetErrorString(cudaStatus) &lt;&lt; " (euclidDist)" &lt;&lt; endl;
          cin.get();
          return 1;
      }
      thrust::copy(thrustC.begin(), thrustC.end(), vectorC.begin());
      cudaStatus = cudaGetLastError();
      if (cudaStatus != cudaSuccess)
      {
          cerr &lt;&lt; "Device to host copy failed: " &lt;&lt; cudaGetErrorString(cudaStatus) &lt;&lt; " (thrustC -&gt; vectorC)" &lt;&lt; endl;
          cin.get();
          return 1;
      }
      end = clock();
      elapsed = double(end - start) / CLOCKS_PER_SEC;
      cout &lt;&lt; "Elapsed time parallel processing is (Thrust): " &lt;&lt; elapsed &lt;&lt; " seconds." &lt;&lt; endl;
  
      cin.get();
      return 0;
  }
  </code></pre>
  
  <p>Suggestions will be appreciated. </p>