title: Why am I not getting I/O-compute overlap with this code?
body: <p>The following program:</p>
  
  <pre><code>#include &lt;iostream&gt;
  #include &lt;array&gt;
  
  using clock_value_t = long long;
  
  __device__ void gpu_sleep(clock_value_t sleep_cycles)
  {
      clock_value_t start = clock64();
      clock_value_t cycles_elapsed;
      do { cycles_elapsed = clock64() - start; }
      while (cycles_elapsed &lt; sleep_cycles);
  }
  
  __global__ void dummy(clock_value_t duration_in_cycles)
  {
      gpu_sleep(duration_in_cycles);
  }
  
  int main()
  {
      const clock_value_t duration_in_clocks = 1e7;
      const size_t buffer_size = 5e7;
      constexpr const auto num_streams = 2;
  
      std::array&lt;char*, num_streams&gt; host_ptrs;
      std::array&lt;char*, num_streams&gt; device_ptrs;
      std::array&lt;cudaStream_t, num_streams&gt; streams;
      for (auto i=0; i&lt;num_streams; i++) {
          cudaMallocHost(&amp;host_ptrs[i], buffer_size);
          cudaMalloc(&amp;device_ptrs[i], buffer_size);
          cudaStreamCreateWithFlags(&amp;streams[i], cudaStreamNonBlocking);
      }
      cudaDeviceSynchronize();
      for (auto i=0; i&lt;num_streams; i++) {
          cudaMemcpyAsync(device_ptrs[i], host_ptrs[i], buffer_size, 
              cudaMemcpyDefault, streams[i]);
          dummy&lt;&lt;&lt;128, 128, 0, streams[i]&gt;&gt;&gt;(duration_in_clocks);
          cudaMemcpyAsync(host_ptrs[i], device_ptrs[i], buffer_size, 
              cudaMemcpyDefault, streams[i]);
      }
      for (auto i=0; i&lt;num_streams; i++) { cudaStreamSynchronize(streams[i]); }
      for (auto i=0; i&lt;num_streams; i++) {
          cudaFreeHost(host_ptrs[i]);
          cudaFree(device_ptrs[i]);
      }
  }
  </code></pre>
  
  <p>should result in overlapping I/O and Compute between the work on the first and second streams: When the first stream's Host-to-Device ends, the first stream's kernel can start, but so can the second stream's Host-to-Device transfer. Instead, I get the following timeline, with no overlap:</p>
  
  <p><a href="https://i.stack.imgur.com/aJRlv.png" rel="nofollow noreferrer"><img src="https://i.stack.imgur.com/aJRlv.png" alt="enter image description here"></a></p>
  
  <p>I <em>think</em> I've covered my bases to ensure overlap. The streams are non-blocking (and indeed the enqueueing of work concludes well before the first HtoD does); the host memory is pinned... so what's missing for me to see overlap?</p>
  
  <p>Using CUDA 8.0.61 on GNU/Linux Mint 18.2 with an NVIDIA GTX 650 Ti Boost. But the driver is v384.59.</p>