TRT2-CUDA_Runtime_Api-2 流

流是一种基于context上的任务管道抽象，一个context可以创建多个流。

流是异步控制的主要方式。

nullptr表示默认流，每个线程都有自己的默认流。

// 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;
    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. stream是一个流句柄，可以当做是一个队列

        - cuda执行器从stream中一条条的读取并执行指令

        - 例如cudaMemcpyAsync函数等同于向stream这个队列中加入一个cudaMemcpy指令并排队

        - 使用到了stream的函数，便立即向stream中加入指令后立即返回，并不会等待指令执行结束

        - 通过cudaStreamSynchronize函数，等待stream中所有指令执行完毕，也就是队列为空

2. 当使用stream时，要注意

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

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

        - `cudaEventCreate`，创建事件

        - `cudaEventRecord`，记录事件，即在stream中加入某个事件，当队列执行到该事件后，修改其状态

        - `cudaEventQuery`，查询事件当前状态

        - `cudaEventElapsedTime`，计算两个事件之间经历的时间间隔，若要统计某些核函数执行时间，请使用这个函数，能够得到最准确的统计

        - `cudaEventSynchronize`，同步某个事件，等待事件到达

        - `cudaStreamWaitEvent`，等待流中的某个事件

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

        - cudaMemcpyAsync(... 默认流) 加入队列

        - cudaStreamSynchronize(默认流) 等待执行完成

        - 默认流与当前设备上下文类似，是与当前设备进行的关联

        - 因此，如果大量使用默认流，会导致性能低下