3. 流同步行为
默认流
默认流,当传入的cudaStream_t参数为0时使用,或由隐式操作流的API调用时,可以配置为具有以下描述的legacy或per-thread同步行为。
可以通过在每个编译单元中使用 --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)来访问传统的默认流。