前言

杜老师推出的 tensorRT从零起步高性能部署 课程，之前有看过一遍，但是没有做笔记，很多东西也忘了。这次重新撸一遍，顺便记记笔记。

本次课程学习精简 CUDA 教程-流的学习，异步任务的管理

课程大纲可看下面的思维导图

1. 流

关于流的知识你需要了解：

1. 流是一种基于 context 之上的任务管道(任务队列)抽象，一个 context 可以创建 n 个流
2. 流是异步控制的主要方式(CUDA 上高性能并发通过流来实现)
3. nullptr 表示默认流，每个线程都是自己的默认流

我们先举个简单的例子来理解串行同步模式，假如你的女朋友想吃东西了，让你去买，那么整个时序图如下所示：

图1-1 串行同步模式时序图

在串行同步模式下女朋友从想吃苹果到吃到苹果这段时间内什么也不能做，需要等待男朋友把苹果买回来了的信息。

你也可以从函数的角度来思考，将发信息买苹果看作一个函数调用，这个函数可能有出门，去买苹果等方法，买回来了之后返回结果，而函数拿到这个结果再做下一步处理，从函数调用到拿到返回值结果这段时间内其实什么也没做，这个就是一个典型的串行同步方式。

我们再来看下异步的方式

图1-2 异步模式时序图

在异步模式下，女朋友发完想吃苹果的消息后就去写作业了，不会傻乎乎的等着，这也符合我们生活的基本行为。可能写了会作业又想吃西瓜了，然后又给男朋友发消息说想吃西瓜，男朋友又屁颠屁颠的去买西瓜，突然可能又想喝奶茶了，然后发消息跟男朋友说想喝奶茶，发完消息后是不是又可以干别的事情，比如打游戏等等。

最后等待男朋友回来即可，当然你可以在任意时候等待拿到你想要的东西，简单来说你可以选择在刚打完游戏的时候就去等(可能你比较渴😂)，你也可以把你的事情忙完后再去等，甚至你可以等到你男朋友买回来一段时间后再去拿东西，这个从什么时间开始去等是你能控制的，也就是说你能决定什么时候去等待拿回你想要得到的结果

从上面的例子中，我们来对比学习下流

1. 上面的例子中，男朋友的微信消息，就是任务队列，流的一种抽象
2. 女朋友发出指令后，她可以做任何事情，无需等待指令执行完毕（指令发出的耗时也是极短的）
3. 也就是说，异步操作，执行的代码加入流的队列后，立即返回，不耽误时间
4. 女朋友发的指令被送到流中排队，男朋友根据流的队列，顺序执行
5. 女朋友可以选择性在需要的时候等待所有的执行结果
6. 新建一个流，就是新建一个男朋友，给他发指令就是给他发微信，你可以新建很多个男朋友
7. 通过 cudaEvent 可以选择性等待任务队列中的部分任务是否就绪

• 比如女朋友在发送买西瓜消息的同时添加了 cudaEvent，也就是等同于告诉男朋友如果你买好了记得给我一个回应，让我知道你已经完成了这个事情

下面我们来看下流案例的示例代码：

// CUDA运行时头文件
#include <cuda_runtime.h>

#include <stdio.h>
#include <string.h>

#define checkRuntime(op)  __check_cuda_runtime((op), #op, __FILE__, __LINE__)

bool __check_cuda_runtime(cudaError_t code, const char* op, const char* file, int line){
    if(code != cudaSuccess){    
        const char* err_name = cudaGetErrorName(code);    
        const char* err_message = cudaGetErrorString(code);  
        printf("runtime error %s:%d  %s failed. \n  code = %s, message = %s\n", file, line, op, err_name, err_message);   
        return false;
    }
    return true;
}

int main(){

    int device_id = 0;
    checkRuntime(cudaSetDevice(device_id));

    cudaStream_t stream = nullptr;
    checkRuntime(cudaStreamCreate(&stream));

    // 在GPU上开辟空间
    float* memory_device = nullptr;
    checkRuntime(cudaMalloc(&memory_device, 100 * sizeof(float)));

    // 在CPU上开辟空间并且放数据进去，将数据复制到GPU
    float* memory_host = new float[100];
    memory_host[2] = 520.25;
    checkRuntime(cudaMemcpyAsync(memory_device, memory_host, sizeof(float) * 100, cudaMemcpyHostToDevice, stream)); // 异步复制操作，主线程不需要等待复制结束才继续

    // 在CPU上开辟pin memory,并将GPU上的数据复制回来 
    float* memory_page_locked = nullptr;
    checkRuntime(cudaMallocHost(&memory_page_locked, 100 * sizeof(float)));
    checkRuntime(cudaMemcpyAsync(memory_page_locked, memory_device, sizeof(float) * 100, cudaMemcpyDeviceToHost, stream)); // 异步复制操作，主线程不需要等待复制结束才继续
    checkRuntime(cudaStreamSynchronize(stream));
    
    printf("%f\n", memory_page_locked[2]);
    
    // 释放内存
    checkRuntime(cudaFreeHost(memory_page_locked));
    checkRuntime(cudaFree(memory_device));
    checkRuntime(cudaStreamDestroy(stream));
    delete [] memory_host;
    return 0;
}

运行效果如下：

图1-3 流案例运行结果

上述代码展示了使用 CUDA 中的流（stream）来进行异步数据传输和内存管理，首先我们通过 cudaStreamCreate 创建了一个流对象，用于管理异步操作。然后使用 cudaMemcpyAsync 将 CPU 上的数据异步复制到 GPU 上，注意 cudaMemcpyAsync 相比于 cudaMemcpy 多了一个参数，也就是流。

在异步复制的时候，发出指令立即返回，并不等待复制完成，因此你可以看到如果你在后面添加打印语句，可以发现数据并没有复制，主线程并不需要等待复制完成，可以继续执行后续操作。

接着同样又利用异步复制将 GPU 上的数据复制到 CPU 上来，最后使用 cudaStreamSynchronize 同步流，确保前面的异步复制操作全部完成(也就是男朋友将东西全买回来了😃)，然后打印获取的结果，那你可以发现真正的耗时其实是发生在这一步的。

通过使用流，可以将数据传输和内存操作于主线程的计算任务异步进行，从而提高了并行性和性能。

对于流的使用，你需要注意的是：

1. 指令发出后，流队列中储存的是指令参数（也就是指针或者形参），不能加入队列后立即释放参数指针，这会导致流队列执行该指令时指针失效而错误
2. 应当在十分肯定流已经不需要这个指针后，才进行修改或者释放，否则会有非预期结果出现

• 比如说当你在执行 cudaMemcpyAsync 后立马执行 delete [] memory_host 将 CPU 上数据释放，那其实复制这个过程是没有完成的，而你又将数据进行释放了，因此会产生一些预期外的结果，这点值得大家注意。因此，你需要确保流已经不需要这个指针后，才对其进行操作
• 举个更简单的例子：比如你给钱让男朋友买西瓜，他没有钱，他刚到店拿好西瓜，你把转的钱撤回去了。那么此时你无法预知他是否会跟店家闹起来矛盾，还是屁颠的回去。如果想得到预期结果，必须得让买完西瓜结束后再处理钱的事情

关于流的知识点需要知道的是：(from chatGPT)

1. stream 是一个流句柄，可以当做是一个队列

• cuda 执行器从 stream 中一条条的读取并执行指令
• 例如 cudaMemcpyAsyn 函数等同于向 stream 这个队列中加入一个 cudaMemcpy 指令并排队
• 使用到了 stream 的函数，便立即向 stream 中加入指令后立即返回，并不会等待指令执行结束
• 通过 cudaStreamSynchronize 函数，等待 stream 中所有指令执行完毕，也就是队列为空

2. 当使用 stream 时，要注意

• 由于异步函数会立即返回，因此传递进入的参数要考虑其生命周期，应确认函数调用结束后再做释放

3. 还可以向 stream 中加入 Event，用以监控是否到达了某个检查点

• cudaEventCreate，创建事件
• cudaEventRecord，记录事件，即在 stream 中加入某个事件，当队列执行到该事件后，修改其状态
• cudaEventQuery，查询事件当前状态
• cudaEventElapsedTime，计算两个事件之前经历的时间间隔，若要统计某些核函数执行时间，请使用这个函数，能够得到最准确的统计
• cudaEventSynchronize，同步某个事件，等待事件到达
• cudaStreamWaitEvent，等待流的某个事件

4. 默认流，对于 cudaMemcpy 等同步函数，其等价于执行了

• cudaMemcpyAsync(… 默认流) 加入队列
• cudaStreamSynchronize(默认流) 等待执行完成
• 默认流与当前设备上下文类似，是与当前设备进行的关联
• 因此，如果大量使用默认流，会导致性能低下

总结

本次课程学习了流，流是 CUDA 编程中异步控制的主要方式。我们通过女朋友让男朋友买东西这个非常生动形象的例子来说明了流的使用，男朋友的微信消息其实就是流的任务队列，而女朋友在发出任务指令后就可以做一些其它的事情，因此可以大大提高效率。

值得注意的是，女朋友可以在需要的时候等待执行的结果，她可以自主决定什么时间点去等待，同时还可以创建多个流来实现并发。

对应到代码层面上，主要通过 cudaMemcpyAsync 这个异步复制函数来讲解流的使用，该函数多了一个流的参数，用来做异步，我们最后通过 cudaStreamSynchronize 函数来统一等待流队列中所有的操作结束，主要的耗时其实也就是在这一步。我们还需要注意的一个点就是流队列中储存的指令参数，你必须确定已经不需要后才能进行修改或者释放，否则会有一些意外的结果。