title: cudamemcpy device->host time changes between iterations
body: <p>I ran into this very weird problem when coding up some CUDA code: the same piece of cudaMemcpy from the gpu to cpu memory takes different time to finish in different iterative calls to the subroutine, and this is a huge difference: ~60 ms vs ~0.02 ms.</p>
  
  <p>The code is as follows:</p>
  
  <pre><code>float calc_formation_obj( int formationNo, bool calcObj )
  {
  
      int i;
  
      int prev = prevCP[aperIndex];
      int next = nextCP[aperIndex];
      float ll = formations_l[formationNo];
      float rl = formations_r[formationNo];
  
      float obj = 0.0;
      float *f_grid = new float[grid_size_voxe];
  
  
  
      // use ll and rl
      thrust::device_ptr&lt;float&gt; dll(d_leafpos_l);
      thrust::device_ptr&lt;float&gt; drl(d_leafpos_r);
  
      dll[rows_per_beam*aperIndex+ rowIndex] = ll;
      drl[rows_per_beam*aperIndex+ rowIndex] = rl;
  
  
      // set all leaf positions between prev/next
      set_leafpos&lt;&lt;&lt;grid_size_ncps,BLOCK_SIZE&gt;&gt;&gt; (aperIndex, rowIndex, prev, next, ncps, d_leafpos_l,
              d_leafpos_r, ll, rl, rows_per_beam, d_cp_angles);
  
  
  
      // copy dose to dose_temp
      thrust::device_ptr&lt;float&gt; ddose(d_dose);
      thrust::device_ptr&lt;float&gt; dtp(d_dose_temp);
  
      thrust::copy(ddose, ddose+nvoxel, dtp);
  
  
      // the angles actually being added
      if (prev==-1) {
          prev = 0;
      }
      if (next==ncps) {
          next = ncps-1;
      }
  
  
  
      // add dose from all these leaf positions
      // if last arg 1 then add
      add_remove_dose&lt;&lt;&lt;grid_size_ncps,BLOCK_SIZE&gt;&gt;&gt; (prev,next, rowIndex, d_dose_temp, d_leafpos_l,
              d_leafpos_r, d_voxe_b, d_dijs_b, d_voxnumperbixcum, d_flu_cp, rows_per_beam, bix_per_row, beamletSize, 1);
  
  
      if (!calcObj) {
          return(0.0);
      }
  
  
      // initialize
      cudaMemset((void*)d_f_voxel, 0, voxesize_f);
      cudaMemset((void*)d_f_grid, 0, sizeof(float)*grid_size_voxe);
  
      // then calculate objective
      calc_obj_dose&lt;&lt;&lt;grid_size_voxe,BLOCK_SIZE&gt;&gt;&gt;( d_dose_temp, d_f_voxel, d_thresh, d_is_target, nvoxel,
              d_f_grid, d_od_wt, d_ud_wt );
  
  
  
      // copy results from GPU
      time_t time_1,time_2;
      float elapse;
      time_1=clock();
      cudaMemcpy(f_grid, d_f_grid, sizeof(float)*grid_size_voxe, cudaMemcpyDeviceToHost);
      time_2 = clock();
      elapse = ((float)time_2 - (float)time_1)/1000;
      printf( "iter %d copy time: %f ms\n", formationNo, elapse );
  
      obj = 0.0;
      for (i=0; i&lt;grid_size_voxe; i++) {
          obj += f_grid[i];
      }
  
  
      delete[] f_grid;
  
      return(obj);
  }
  </code></pre>
  
  <p>This subroutine is called many times during the program, and everytime it runs I record the runtime for </p>
  
  <pre><code>cudaMemcpy(f_grid, d_f_grid, sizeof(float)*grid_size_voxe, cudaMemcpyDeviceToHost);
  </code></pre>
  
  <p>and the results I got look like:</p>
  
  <pre><code>iter 0 copy time: 0.018000 ms
  iter 1 copy time: 66.445999 ms
  iter 2 copy time: 64.239998 ms
  iter 3 copy time: 66.959999 ms
  iter 4 copy time: 66.328003 ms
  iter 5 copy time: 65.656998 ms
  iter 6 copy time: 66.120003 ms
  iter 7 copy time: 63.811001 ms
  iter 8 copy time: 66.530998 ms
  iter 9 copy time: 65.686996 ms
  iter 10 copy time: 65.808998 ms
  iter 11 copy time: 0.027000 ms
  iter 12 copy time: 64.346001 ms
  iter 13 copy time: 66.407997 ms
  iter 14 copy time: 65.796997 ms
  iter 15 copy time: 65.471001 ms
  iter 16 copy time: 66.209000 ms
  iter 17 copy time: 63.799000 ms
  iter 18 copy time: 66.542999 ms
  iter 19 copy time: 65.660004 ms
  iter 20 copy time: 65.102997 ms
  iter 21 copy time: 0.019000 ms
  iter 22 copy time: 64.665001 ms
  iter 23 copy time: 66.653999 ms
  iter 24 copy time: 65.648003 ms
  iter 25 copy time: 65.233002 ms
  iter 26 copy time: 65.851997 ms
  iter 27 copy time: 63.992001 ms
  iter 28 copy time: 66.172997 ms
  iter 29 copy time: 65.503998 ms
  iter 30 copy time: 0.020000 ms
  iter 31 copy time: 66.277000 ms
  iter 32 copy time: 63.881001 ms
  iter 33 copy time: 66.537003 ms
  iter 34 copy time: 65.626999 ms
  iter 35 copy time: 65.387001 ms
  iter 36 copy time: 66.084999 ms
  iter 37 copy time: 63.797001 ms
  iter 38 copy time: 0.017000 ms
  iter 39 copy time: 65.707001 ms
  iter 40 copy time: 65.553001 ms
  iter 41 copy time: 66.362999 ms
  iter 42 copy time: 63.634998 ms
  </code></pre>
  
  <p>This is run on Mac OSX with a GeForce GTX 285, using CUDA 4.2. I've no idea why it's doing this --- should be a straightforward copy. Any help is appreciated!!</p>