cuda中的两种默认stream

背景

一般大家写cuda代码用到流的情况不多，很多使用简单的使用默认流就行，但是对于某些应用使用多流比较合适：

1. 需要平凡调用很多kernel，但是每个kernel只能使用一点gpu资源的时候
2. 当一些数据copy可以和计算重叠时候
   对于kernel并行的基础知识可以参考这里, 在文章的结尾简单解释一下默认流的问题，当时为了避免同步的影响，引入了cudaStreamNonBlocking，最近在看一些框架发现原来有更好的解决方式，下面对这个知识点进行一个简单的总结。

默认流

nvidia的默认流有两种：

• legacy(传统的或者停产的)
• per-thread

legacy

legacy是一个隐式流(也就说kernel launch的时候不用写)会自动同步一个CUcontext（如果只是使用runtime api的话，每个设备会有一个context）中的所有非non-blocking流，显示用的话可以把cudaStreamLegacy传给kernel调用。

demo

const int N = 1 << 20;

__global__ void kernel(float *x, int n)
{
   
   
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    for (int i = tid; i < n; i += blockDim.x * gridDim.x) {
   
   
        x[i] = sqrt(pow(3.14159,i));
    }
}

int main()
{
   
   
    const int num_streams = 8;

    cudaStream_t streams[num_streams];
    float *data[num_streams];

    for (int i = 0; i < num_streams; i++) {
   
   
        cudaStreamCreate(&streams[i]);
 
        cudaMalloc(&data[i