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) 显式使用。