title: Mixing implicit and explicit CUDA streams
body: <p>I have some CUDA 8.0 code (edit: that I inherited, not something I wrote) that looks basically like this:</p>
  <pre><code>cudaMemcpy(devInputData, ..., cudaMemcpyHostToDevice);
  kernelThings&lt;&lt;&lt;GRIDS, BLOCKS, 0, myStream&gt;&gt;&gt;(devInputData);
  cudaDeviceSynchronize();
  </code></pre>
  <p><code>cudaMemcpy()</code> works synchronously without a stream, so this code is correct, as far as I understand.</p>
  <p>If I compile using <a href="https://docs.nvidia.com/cuda/cuda-runtime-api/stream-sync-behavior.html" rel="nofollow noreferrer"><code>CUDA_API_PER_THREAD_DEFAULT_STREAM</code></a>, is this code still safe? I argue no, the <code>cudaMemcpy()</code> now happens asynchronously, so there is a potential for the kernel to start before <code>cudaMemcpy()</code> finishes. However, looking at the <code>Nsight</code> profiler, I see there is no overlap - textually I saw:</p>
  <pre><code>[Memcpy HtoD]
                    [kernelThings]
  </code></pre>
  <p>With a 16 microsecond gap between the two functions. This behaviour was repeated multiple times within the application.</p>
  <p>However, I next deleted <code>cudaDeviceSynchronize()</code>, reran <code>Nsight</code>, and saw that now they overlapped:</p>
  <pre><code>[Memcpy HtoD]
           [kernelThings]
  </code></pre>
  <p>The kernel now started 10 microseconds before the <code>cudaMemcpy</code> finished.</p>
  <p>Obviously, the correct fix is to use a stream with <code>cudaMemcpyAsync()</code> too:</p>
  <pre><code>cudaMemcpyAsync(devInputData, ..., cudaMemcpyHostToDevice, myStream);
  </code></pre>
  <p>However, my question is why did I not see the processes overlap when using <code>cudaDeviceSynchronize()</code>? Is the simple answer that I might not be so lucky with a different version of CUDA or GPU?</p>