title: CUDA speedup for simple calculations
body: <p>I have the following code in <code>cuda_computation.cu</code></p>
  
  <pre><code>#include &lt;iostream&gt;
  #include &lt;stdio.h&gt;
  #include &lt;cuda.h&gt;
  #include &lt;assert.h&gt;
  
  void checkCUDAError(const char *msg);
  
  __global__ void euclid_kernel(float *x, float* y, float* f)
  {
    int idx = blockIdx.x*blockDim.x + threadIdx.x;
    int i = blockIdx.x;
    int j = threadIdx.x;
    f[idx] = sqrt((x[i]-x[j])*(x[i]-x[j]) + (y[i]-y[j])*(y[i]-y[j]));
  }
  int main()
  {
    float *xh;
    float *yh;
    float *fh;
    float *xd;
    float *yd;
    float *fd;
  
    size_t n = 256;
    size_t numBlocks = n;
    size_t numThreadsPerBlock = n;
  
    size_t memSize = numBlocks * numThreadsPerBlock * sizeof(float);
    xh = (float *) malloc(n * sizeof(float));
    yh = (float *) malloc(n * sizeof(float));
    fh = (float *) malloc(memSize);
  
    for(int ii(0); ii!=n; ++ii)
      {
        xh[ii] = ii;
        yh[ii] = ii;
      }
  
    cudaMalloc( (void **) &amp;xd, n * sizeof(float) );
    cudaMalloc( (void **) &amp;yd, n * sizeof(float) );
    cudaMalloc( (void **) &amp;fd, memSize );
    for(int run(0); run!=10000; ++run)
      {
        //change value to avoid optimizations
        xh[0] = ((float)run)/10000.0;
        cudaMemcpy( xd, xh, n * sizeof(float), cudaMemcpyHostToDevice );
        checkCUDAError("cudaMemcpy");
        cudaMemcpy( yd, yh, n * sizeof(float), cudaMemcpyHostToDevice );
        checkCUDAError("cudaMemcpy");
        dim3 dimGrid(numBlocks);
        dim3 dimBlock(numThreadsPerBlock);
        euclid_kernel&lt;&lt;&lt; dimGrid, dimBlock &gt;&gt;&gt;( xd, yd, fd );
        cudaThreadSynchronize();
        checkCUDAError("kernel execution");
        cudaMemcpy( fh, fd, memSize, cudaMemcpyDeviceToHost );
        checkCUDAError("cudaMemcpy");
      }
    cudaFree(xd);
    cudaFree(yd);
    cudaFree(fd);
    free(xh);
    free(yh);
    free(fh);
    return 0;
  }
  
  void checkCUDAError(const char *msg)
  {
    cudaError_t err = cudaGetLastError();
    if( cudaSuccess != err) 
      {
        fprintf(stderr, "Cuda error: %s: %s.\n", msg, cudaGetErrorString( err) );
        exit(-1);
      }                         
  }
  </code></pre>
  
  <p>It takes about 6" to run on an FX QUADRO 380, while the corresponding serial version using just one i7-870 core takes just about 3". Do I miss something? Is the code under optimised in some ways? Or is it just expected behaviour that for simple calculations (like this all-pairs Euclidean distance) the overhead needed to move memory exceeds the computational gain?</p>