title: CUDA default streams and CUDA_API_PER_THREAD_DEFAULT_STREAM
body: <p><a href="https://docs.nvidia.com/cuda/cuda-runtime-api/stream-sync-behavior.html" rel="nofollow noreferrer">The documentation here</a> tries to explain how default streams are handled.</p>
  
  <p>Given code like this (ignoring allocation errors):</p>
  
  <pre><code>char *ptr;
  char source[1000000];
  cudaMalloc((void**)&amp;ptr, 1000000);
  cudaMemcpyAsync(ptr, source, 1000000, cudaMemcpyHostToDevice);
  myKernel&lt;&lt;&lt;1000, 1000&gt;&gt;&gt;(ptr);
  </code></pre>
  
  <p>Is there a risk that <code>myKernel</code> will start before <code>cudaMemcpyAsync</code> finishes copying? I think "No" because this is a "Legacy default stream" as described in the documentation.</p>
  
  <p>However, if I compile with <code>CUDA_API_PER_THREAD_DEFAULT_STREAM</code> what happens? The text for "Per-thread default stream" says:</p>
  
  <blockquote>
    <p>The per-thread default stream is an implicit stream local to both the thread and the <code>CUcontext</code>, and which does not synchronize with other streams (just like explcitly created streams). The per-thread default stream is not a non-blocking stream and will synchronize with the legacy default stream if both are used in a program.</p>
  </blockquote>
  
  <p>I think this might also be OK as both <code>cudaMemcpyAsync</code> and <code>myKernel</code> are effectively using <code>CU_STREAM_PER_THREAD</code>; am I correct?</p>
  
  <hr>
  
  <p><sub>The reason I ask is that I have a really weird intermittent CUDA error 77 in a kernel that I can only explain by a <code>cudaMemcpyAsync</code> not finishing before calling <code>myKernel</code>, which would mean that I am not understanding the documentation. The real code is too involved and too proprietary to make an MCVE, though.</sub></p>