title: Why is the cuda version 2x slower than the cpu code?
body: <p>This is a little baffling to me why the cuda code runs about twice as slow as the cpu version.  The cpu code is commented out above the main.  I am just counting all the primes from 0 to (512 * 512 * 512).  The cpu version executed in about 97 seconds whereas the gpu version took 182 seconds. I have an intel i7 running at 4 Ghz and an nvidia GTX 960.  Any ideas why?</p>
  
  <pre><code>#include &lt;cuda.h&gt;
  #include &lt;iostream&gt;
  #include &lt;cstdint&gt;
  #include &lt;stdio.h&gt;
  #include &lt;ctime&gt;
  #include &lt;vector&gt;
  #include &lt;cstdlib&gt;
  #include &lt;climits&gt;
  
  using namespace std;
  
  
  
  __host__ __device__  bool is_prime(uint32_t n)
  {
      if(n == 2)
          return true;
      if(n % 2 == 0)
          return false;
      uint32_t sr = sqrtf(n);
  
      for(uint32_t i = 3; i &lt;= sr; i += 2)
          if(n % i == 0)
              return false;
      return true;
  }
  
  __global__ void prime_sum(unsigned int* count)
  {
      uint32_t n = (blockIdx.y * gridDim.y + blockIdx.x) * blockDim.x + threadIdx.x;
      if(is_prime(n))
          atomicAdd(count, 1);
  }
  
  int main()
  {
      /* CPU VERSION
      time_t start = time(0);
      int pcount = 0;
      for(uint32_t i = 0; i &lt; (512 * 512 * 512); i++)
      {
          if(is_prime(i)) pcount++;
      }
      start = time(0) - start;
      std::cout &lt;&lt; pcount &lt;&lt; "\t" &lt;&lt; start &lt;&lt; std::endl;
      */
  
      //CUDA VERSION
      time_t start = time(0);
      unsigned int* sum_d;
      cudaMalloc(&amp;sum_d, sizeof(unsigned int));
      cudaMemset(sum_d, 0, sizeof(unsigned int));
  
      prime_sum&lt;&lt;&lt; dim3(512, 512), 512 &gt;&gt;&gt;(sum_d);
  
      unsigned int sum = 0;
      cudaMemcpy(&amp;sum, sum_d, sizeof(unsigned int), cudaMemcpyDeviceToHost);
      start = time(0) - start;
      std::cout &lt;&lt; sum &lt;&lt; "\t" &lt;&lt; start &lt;&lt; std::endl;
      cudaFree(sum_d);
  
      return 0;
  }
  </code></pre>
  
  <p>Here is one idea.  The efficiency of the is_prime function comes from being able to exit quickly most of the time because most numbers will be divisible by 2 or lower numbers so when executed in serial most of the time the loop exits fast.  However due to warps each group of 32 threads must wait for the worst to finish.  Also I am including evens so half the threads will be eliminated by the first if.</p>