CUDA默认流的同步行为

默认流
对于需要指定 cudaStream_t参数的 cuda API，如果将 0作为实参传入，则视为使用默认流；对于不需要指定 cudaStream_t参数的 cuda API，则也视为使用默认流。
在 cuda中，默认流有两种类型，一种是 legacy默认流，一种是 per-thread默认流，这两种默认流的同步行为不一样，在使用的时候需要注意一下。具体使用哪种默认流，有以下三种方式进行指定：

1. 不指定，默认使用 legacy默认流；
2. 在编译 cuda程序时，通过 nvcc的 --default-stream进行指定，可选的取值是：{legacy|per-thread}；
3. 在需要传入 cudaStream_t/ CUstream参数的时候，指定：cudaStreamLegacy/ CU_STREAM_LEGACY（legacy默认流），或 cudaStreamPerThread/ CU_STREAM_PER_THREAD（per-thread默认流）。
   下面重点介绍一下 legacy默认流和 per-thread默认流在同步行为上的差异。

legacy 默认流
legacy默认流会与同一个 CUcontext（如果是使用 runtime的 API，则每个设备对应一个 CUcontext）中的其他流都进行同步，但 non-blocking的流除外。也即是，在执行 legacy默认流中的任务之前，会先等待其他所有 blocking的流执行完成，然后开始执行 legacy默认流中的任务，并且在 legacy默认流后面出现的其他 blocking流中的任务，会先等待 legacy默认流中的任务执行完成，再开始执行。

假设，下面的代码在流 s中 launch了核函数 k_1，在 legacy默认流中 launch了核函数 k_2，在流 s中 launch了核函数 k_3：

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

则上述代码的同步行为是 k_2会被 k_1阻塞，k_3会被 k_2阻塞。

但如果是 non-blocking流，则不会出现上述同步行为，也即是，下面三个核函数会存在并行执行的情况，例如：

cudaStream_t s;
cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking);
k_1<<<1, 1, 0, s>>>();
k_2<<<1, 1>>>();
k_3<<<1, 1, 0, s>>>();

per-thread 默认流
在一个 CUcontext内，每个线程都有一个 per-thread默认流，这个流不与其他流进行同步（就像是一个显式创建的流那样），但如果在一个程序中同时使用了 legacy默认流和 per-thread默认流，则 per-thread默认流会与 legacy默认流保持同步。