关于CUDA流[1]
CUDA流表示一个GPU操作队列，该队列中的操作将以添加到流中的先后顺序而依次执行。可以将一个流看做是GPU上的一个任务，不同任务可以并行执行。
CUDA流同步[2]
- 默认流是一个隐式流（无需显式创建，CUDA中默认存在），它与同一
CUcontext[3]中的所有其他流同步，非阻塞流（non-blocking streams）除外。
例如，在如下代码中，在流 s 中启动核函数 k_A，然后在默认流中启动 k_B，在流 s 中启动 k_C。
k_A<<<1, 1, 0, s>>>();
k_B<<<1, 1>>>();
k_C<<<1, 1, 0, s>>>();
其运行结果k_A阻塞k_B的运行，k_B阻塞k_C的运行。
- 如果操作[4]A和操作B位于同一个CUDA流中，操作A和操作B按照先后顺序执行；
例如，在如下代码中，在流 s 中先后启动核函数 k_A和 k_B。
k_A<<<1, 1, 0, s>>>();
k_B<<<1, 1, 0, s>>>();
其运行结果k_A阻塞k_B的运行。
- 如果两个操作位于不同的CUDA流中，则操作A和操作B并没有确定的执行顺序；
例如，在如下代码中，先在流 s1 中先后启动核函数 k_A1和 k_B1，然后在流 s2 中先后启动核函数 k_A2和 k_B2。
k_A1<<<1, 1, 0, s1>>>();
k_B1<<<1, 1, 0, s1>>>();
k_A2<<<1, 1, 0, s2>>>();
k_B2<<<1, 1, 0, s2>>>();
在实际运行结果中，流s1的操作并不会影响流s2的操作，k_A1会阻塞k_B1的运行，而不会阻塞k_B2的运行。
CUDA流管理[5]
创建流
在CUDA中，流的创建有三种方式：cudaStreamCreate，cudaStreamCreateWithFlags，cudaStreamCreateWithPriority
创建CUDA流
__host__ cudaError_t cudaStreamCreate ( cudaStream_t* pStream )
创建不同标记的CUDA流
__host__ cudaError_t cudaStreamCreateWithFlags ( cudaStream_t* pStream , unsigned int flags )
-
cudaStreamDefault: CUDA流的默认创建标记； -
cudaStreamNonBlocking: 创建非阻塞流。在通过该标记创建的CUDA流中运行的工作可以与流0（NULL 流）中的工作同时运行，并且通过该标记创建的流不应与流0执行隐式同步（implicit synchronization）。
创建不同优先级的CUDA流
__host__ cudaError_t cudaStreamCreateWithPriority ( cudaStream_t* pStream , unsigned int flags, int priority )
通过priority参数，指定CUDA 流的优先级，数值越小其优先级越高（“0”表示默认优先级）。高优先级流中的计算核函数（compute kernels）可以抢占已经在低优先级流中执行的核函数，而主机和设备之间的内存操作不受优先级的影响。
销毁流
__host____device__ cudaError_t cudaStreamDestroy ( cudaStream_t stream )
销毁并清理异步流（asynchronous stream）。
流同步操作
__host__ cudaError_t cudaStreamSynchronize ( cudaStream_t stream )
调用该操作时会发生阻塞，等待指定流中操作完成。
-
这里的操作通常指异步内存拷贝（cudaMemcpyAsync）和核函数调用。 ↩
