title: Thrust Complex Inner Product Run on GPU Slower Than STL Implementation on CPU
body: <p>I've got the following two implementations of computing a complex inner product, one using STL libraries running on the CPU and one using Thrust running on the GPU:</p>
  
  <p><b>CPU Implementation</b></p>
  
  <pre><code>#include &lt;vector&gt;
  #include &lt;numeric&gt;
  #include &lt;complex&gt;
  
  int main(int argc, char **argv)
  {
      int vec_size = atoi(argv[1]);
  
      std::vector&lt; std::complex&lt;float&gt; &gt; host_x( vec_size );
      std::generate(host_x.begin(), host_x.end(), std::rand);
  
      std::vector&lt; std::complex&lt;float&gt; &gt; host_y( vec_size );
      std::generate(host_y.begin(), host_y.end(), std::rand);
  
      std::complex&lt;float&gt; z = std::inner_product(host_x.begin(), host_x.end(), host_y.begin(), std::complex&lt;float&gt;(0.0f,0.0f) );
  
      return 0;
  }
  </code></pre>
  
  <p><b>GPU Implementation</b></p>
  
  <pre><code>#include &lt;thrust/host_vector.h&gt;
  #include &lt;thrust/device_vector.h&gt;
  #include &lt;thrust/inner_product.h&gt;
  #include &lt;thrust/complex.h&gt;
  
  int main(int argc, char **argv)
  {
      int vec_size = atoi(argv[1]);
  
      thrust::host_vector&lt; thrust::complex&lt;float&gt; &gt; host_x( vec_size );
      thrust::generate(host_x.begin(), host_x.end(), rand);
  
      thrust::host_vector&lt; thrust::complex&lt;float&gt; &gt; host_y( vec_size );
      thrust::generate(host_y.begin(), host_y.end(), rand);
  
      thrust::device_vector&lt; thrust::complex&lt;float&gt; &gt; device_x = host_x;
      thrust::device_vector&lt; thrust::complex&lt;float&gt; &gt; device_y = host_y;
  
      thrust::complex&lt;float&gt; z = thrust::inner_product(device_x.begin(), device_x.end(), device_y.begin(), thrust::complex&lt;float&gt;(0.0f,0.0f) );
  
      return 0;
  }
  </code></pre>
  
  <p>I'm compiling the CPU implementation using g++ and the GPU implementation using mvcc. Both have -O3 optimizations on. I run both implementations with 3,000,000 elements in the vector and get the following timing results:</p>
  
  <p>CPU:<br>
  real    0m0.159s<br>
  user    0m0.100s<br>
  sys 0m0.048s<br></p>
  
  <p>GPU:<br>
  real    0m0.284s<br>
  user    0m0.190s<br>
  sys 0m0.083s<br></p>
  
  <p>I'm using the following pieces of software:</p>
  
  <pre><code>$ gcc -v
  Configured with: --prefix=/Applications/Xcode.app/Contents/Developer/usr --with-gxx-include-dir=/Applications/Xcode.app/Contents/Developer/Platforms/MacOSX.platform/Developer/SDKs/MacOSX10.9.sdk/usr/include/c++/4.2.1
  Apple LLVM version 5.1 (clang-503.0.40) (based on LLVM 3.4svn)
  Target: x86_64-apple-darwin13.3.0
  Thread model: posix
  
  $ nvcc --version
  nvcc: NVIDIA (R) Cuda compiler driver
  Copyright (c) 2005-2013 NVIDIA Corporation
  Built on Thu_Sep__5_10:17:14_PDT_2013
  Cuda compilation tools, release 5.5, V5.5.0
  </code></pre>
  
  <p>Along with the latest version of Thrust from the GitHub repo.</p>
  
  <p>My CPU is a 2.4 GHz Intel Core 2 Duo and my GPU is a NVIDIA GeForce 320M 256 MB.</p>
  
  <p><b>Question:</b>
  I'm new to the use of Thrust, but shouldn't my GPU implementation be significantly faster than my CPU implementation? I realize that there are memory transaction costs with GPUs, but I guess I'm trying to figure out if I'm using Thrust correctly to execute the inner product on the GPU since the timing results are unexpectedly reversed in my opinion.</p>
  
  <p><b>EDIT:</b>
  Per everyone's suggestions I made the number of iterations configurable and changed the granularity of the timing as follows:</p>
  
  <pre><code>#include &lt;stdio.h&gt;
  
  #include &lt;thrust/host_vector.h&gt;
  #include &lt;thrust/device_vector.h&gt;
  #include &lt;thrust/inner_product.h&gt;
  #include &lt;thrust/complex.h&gt;
  #include &lt;thrust/execution_policy.h&gt;
  
  int main(int argc, char **argv)
  {
      int vec_size = atoi(argv[1]);
      int iterations = atoi(argv[2]);
  
      float milliseconds = 0;
  
      cudaEvent_t start, stop;
      cudaEventCreate(&amp;start);
      cudaEventCreate(&amp;stop);
  
      thrust::host_vector&lt; thrust::complex&lt;float&gt; &gt; host_x( vec_size );
      thrust::generate(host_x.begin(), host_x.end(), rand);
  
      thrust::host_vector&lt; thrust::complex&lt;float&gt; &gt; host_y( vec_size );
      thrust::generate(host_y.begin(), host_y.end(), rand);
  
      printf("vector size = %lu bytes\n", vec_size * sizeof(thrust::complex&lt;float&gt;)); 
  
      cudaEventRecord(start);
  
      thrust::device_vector&lt; thrust::complex&lt;float&gt; &gt; device_x = host_x;
      thrust::device_vector&lt; thrust::complex&lt;float&gt; &gt; device_y = host_y;
      cudaEventRecord(stop);
  
      cudaEventSynchronize(stop);
      cudaEventElapsedTime(&amp;milliseconds, start, stop);
  
      printf("copy (device)\t\t%f ms\n", milliseconds);
  
      cudaEventRecord(start);
  
      for(int i = 0; i &lt; iterations; ++i)
      {
          thrust::inner_product(thrust::cuda::par, device_x.begin(), device_x.end(), device_y.begin(), thrust::complex&lt;float&gt;(0.0f,0.0f) );
      }
  
      cudaEventRecord(stop);
  
      cudaEventSynchronize(stop);
      cudaEventElapsedTime(&amp;milliseconds, start, stop);
  
      printf("inner_product (device)\t%f ms\n", milliseconds/iterations); 
  
      cudaEventRecord(start);
  
      for(int i = 0; i &lt; iterations; ++i)
      {
          thrust::inner_product(thrust::host, host_x.begin(), host_x.end(), host_y.begin(), thrust::complex&lt;float&gt;(0.0f,0.0f) );
      }
  
      cudaEventRecord(stop);
  
      cudaEventSynchronize(stop);
      cudaEventElapsedTime(&amp;milliseconds, start, stop);
  
      printf("inner_product (host)\t%f ms\n", milliseconds/iterations);   
  
      return 0;
  }
  </code></pre>
  
  <p>On a Tegra K1 I got the following:</p>
  
  <pre><code>$ nvcc complex_inner_product.cu -O3 -arch=sm_32 -o cip
  $ ./cip 3100000 1000
  vector size = 24800000 bytes
  copy (device)       45.741653 ms
  inner_product (device)  10.595121 ms
  inner_product (host)    1.807912 ms
  </code></pre>
  
  <p>On an Intel Core 2 Duo 2.4 GHz and GeForce 320M I got the following results:</p>
  
  <pre><code>$ nvcc complex_inner_product.cu -O3 -arch=sm_12 -o cip
  $ ./cip 3100000 1000
  vector size = 24800000 bytes
  copy (device)       227.765213 ms
  inner_product (device)  42.180416 ms
  inner_product (host)    0.000018 ms
  </code></pre>
  
  <p>On an Intel Core i5 3.3 GHz and GeForce GT 755M:</p>
  
  <pre><code>$ nvcc complex_inner_product.cu -O3 -arch=sm_30 -o cip
  $ ./cip 3100000 1000
  vector size = 24800000 bytes
  copy (device)       22.930016 ms
  inner_product (device)  6.249663 ms
  inner_product (host)    0.000003 ms
  </code></pre>
  
  <p>So no matter what compute capability or hardware I use, the host processor is at least 10x faster than the GPU. Any ideas?</p>