cuda中的两种默认stream

文章目录

• • 背景
  • 默认流
  • legacy
  • • demo
  • per-thread
  • • demo
  • 扩展
  • `注意`

背景

一般大家写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], N * sizeof(float));
        
        // launch one worker kernel per stream
        kernel<<<1, 64, 0, streams[i]>>>(data[i], N);

        // launch a dummy kernel on the default stream
        kernel<<<1, 1>>>(0, 0);
    }

    cudaDeviceReset();

    return 0;
}

接下来进行编译：nvcc ./stream_test.cu -o stream_legacy， 可以得到下面的结果，可以看到其他流会被legacy默认流给同步。

per-thread

per-thread也是一个隐式流，和legacy不同的是他不会同步其他stream，如果per-thread和legacy同时出现的话会和leagcy同步（没懂啥意思）。

• 显式使用的话就是cudaStreamPerThread。
• 隐式使用的话，由于在cuda的使用过程中默认是legacy，如果想用per-thread也需要编译器参与，可以给nvvc传入参数–default-stream per-thread, 或者定义一个宏CUDA_API_PER_THREAD_DEFAULT_STREAM，在include任何CUDA headers之前，也是就说放在cmake里面。

demo

还是上面的代码，我们先看看如何启动per-thread。

• 隐式编译：nvcc --default-stream per-thread ./stream_test.cu -o stream_per-thread
  得到的结果如下：
  
  可以看到确实没有再去同步其他tream了。
• 隐式定义宏：nvcc -DCUDA_API_PER_THREAD_DEFAULT_STREAM ./stream_test.cu -o stream_thread ，可以得到同样的结果。这里注意不要把宏定义在文件里，这样没啥用
• 显示使用cudaStreamPerThread，这里需要修改一点代码

kernel<<<1, 1,0, cudaStreamPerThread>>>(0, 0);

进行编译：nvcc ./stream_test.cu -o stream_thread_explict可以得到下面的结果：

扩展

其实per-thread并不是为了取代legacy的目的而设计的，只是恰巧有no-blocking stream的效果而已。本质上per-thread是为了多线程设计的流模式，当面对多线程的进程应用的时候，如果每个线程都使用的是legacy stream的时候，那么多线程的意义其实就失去了，因为在gpu上还是串行的。

#include 
#include 

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));
    }
}

void *launch_kernel(void *dummy)
{
    float *data;
    cudaMalloc(&data, N * sizeof(float));

    kernel<<<1, 64>>>(data, N);

    cudaStreamSynchronize(0);

    return NULL;
}

int main()
{
    const int num_threads = 8;

    pthread_t threads[num_threads];

    for (int i = 0; i < num_threads; i++) {
        if (pthread_create(&threads[i], NULL, launch_kernel, 0)) {
            fprintf(stderr, "Error creating threadn");
            return 1;
        }
    }

    for (int i = 0; i < num_threads; i++) {
        if(pthread_join(threads[i], NULL)) {
            fprintf(stderr, "Error joining threadn");
            return 2;
        }
    }

    cudaDeviceReset();

    return 0;
}

我们首先使用legacy看看会是什么效果， 编译命令：nvcc ./pthread_test.cu -o pthreads_legacy

可以看到只有一个default , 接下来看看用per-thread, 编译命令：nvcc --default-stream per-thread ./pthread_test.cu -o pthreads_per_thread 结果如下，会8个stream并行：

再来看看看legacy和per-threadt同时出现在一个进程中会出现什么情况，首先先修改一下代码：

__global__ void defalut_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));
    }
}

void *launch_kernel(void *dummy)
{
    float *data;
    cudaMalloc(&data, N * sizeof(float));

    kernel<<<1, 64>>>(data, N);
    defalut_kernel<<<1, 64, 0, cudaStreamLegacy>>>(data, N);

    cudaStreamSynchronize(0);

    return NULL;
}

进行编译nvcc --default-stream per-thread ./pthread_test.cu -o pthreads_per_thread 可以得到如下结果，可以看到legacy就一个stream, 但是依然把整个进程中的所有流给同步了。

注意

cudaMemcpy理论也使用一个默认流，在两种不同的启用per-thread模式时候，效果完全不一样：

• 编译时传–default-stream per-thread， 如果cudaMemcpy放在子线程中，那么cudaMemcpy使用的就是per-thread(这里在nsys上是看不到default stream，也就是普通的per-thread stream), 而如果这个函数用在的主线程中，那么就会使用legacy stream, 在nsys上可以看到default stream，即使传入–default-stream per-thread也不会影响主线程中的默认流时legacy stream。（好像有随机性，有时候复现不出来）
• 编译时候传入宏-DCUDA_API_PER_THREAD_DEFAULT_STREAM， 如果cudaMemcpy放在子线程中，那么cudaMemcpy使用的就是per-thread， 如果在主进程里，那么会使用的是per-thread stream(也就是新建一个stream)。