3. 流同步行为

默认流

当 0 作为 cudaStream_t 传递或由隐式操作流的 API 使用时，默认流可以配置为具有旧版或每线程同步行为，如下所述。

行为可以通过 --default-stream nvcc 选项按编译单元进行控制。或者，可以通过在包含任何 CUDA 头文件之前定义 CUDA_API_PER_THREAD_DEFAULT_STREAM 宏来启用每线程行为。无论哪种方式，CUDA_API_PER_THREAD_DEFAULT_STREAM 宏都将在使用每线程同步行为的编译单元中定义。

旧版默认流

旧版默认流是一个隐式流，它与同一 CUcontext 中的所有其他流同步，除了下面描述的非阻塞流。（对于仅使用运行时 API 的应用程序，每个设备将有一个上下文。）当在旧版流中执行操作（例如内核启动或 cudaStreamWaitEvent()）时，旧版流首先等待所有阻塞流，操作在旧版流中排队，然后所有阻塞流等待旧版流。

例如，以下代码在流 s 中启动内核 k_1，然后在旧版流中启动 k_2，然后在流 s 中启动 k_3

k_1<<<1, 1, 0, s>>>();
k_2<<<1, 1>>>();
k_3<<<1, 1, 0, s>>>();

结果行为是 k_2 将阻塞 k_1，而 k_3 将阻塞 k_2。

可以使用带有流创建 API 的 cudaStreamNonBlocking 标志创建不与旧版流同步的非阻塞流。

旧版默认流可以与 CUstream (cudaStream_t) 句柄 CU_STREAM_LEGACY (cudaStreamLegacy) 显式使用。

每线程默认流

每线程默认流是一个隐式流，它在线程和 CUcontext 本地，并且不与其他流同步（就像显式创建的流一样）。每线程默认流不是非阻塞流，如果程序中同时使用两者，它将与旧版默认流同步。

每线程默认流可以与 CUstream (cudaStream_t) 句柄 CU_STREAM_PER_THREAD (cudaStreamPerThread) 显式使用。