cuda流stream，异步任务的管理

句柄： 句柄可以认为是系统对资源（如线程）的分配的一个编号，是一个long类型的数字，关闭这个编号，对于不同的资源，效果不尽相同。对于线程来说，关闭这个编号并不意味着终止线程，只是之后很难再操纵这个线程。
线程、线程句柄、线程ID

1.stream是一个流句柄，可以当做是一个队列，流是异步控制的主要方式（异步操作，发出指令立即返回，即立马执行下一行代码，而同步操作是要等待操作执行结束后才返回）

流是一种基于context之上的任务的队列，一个context可以创建n个流，nullptr表示默认流，每个线程都有自己的默认流。
cuda执行器从stream中一条条的读取并执行指令
例如cudaMemcpyAsync函数等同于向stream这个队列中加入一个cudaMemcpy指令并排队
使用到了stream的函数，便立即向stream中加入指令后立即返回，并不会等待指令执行结束
通过cudaStreamSynchronize函数，等待stream中所有指令执行完毕，也就是队列为空

2.当使用stream时，要注意

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

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

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

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

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

5.代码示例

// CUDA运行时头文件
#include 

#include 
#include 

#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;

    //带Async为异步函数，异步操作，发出指令立即返回，即立马执行下一行，而同步操作是要等待操作执行结束后才返回
    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
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53