（番外）如何将cuda数据存入std::queue实现异步效果

1、std::queue列队如何实现异步？

std::queue是一种队列容器，对于需要先进先出的任务非常实用。比如下面是一个列队的使用用例：

#include <iostream>
#include <queue>

int main() {
    // 创建一个 int 类型的队列
    std::queue<int> myQueue;

    // 向队列中添加一些元素，此时myQueue中存放的数据为->[30,20,10]->
    myQueue.push(10);
    myQueue.push(20);
    myQueue.push(30);

    // 访问队列的头部元素
    std::cout << "Front of the queue: " << myQueue.front() << std::endl;

    // 删除队列的头部元素，此时myQueue中存放的数据为->[30,20]->
    myQueue.pop();

    // 再次访问队列的头部元素，此时myQueue中存放的数据为->[30,20]->
    std::cout << "Front of the queue after pop: " << myQueue.front() << std::endl;

    // 检查队列是否为空，此时myQueue中存放的数据为->[30,20]->
    if (myQueue.empty()) {
        std::cout << "Queue is empty." << std::endl;
    } else {
        std::cout << "Queue is not empty." << std::endl;
    }

    // 获取队列的大小，此时myQueue中存放的数据为->[30,20]->
    std::cout << "Size of the queue: " << myQueue.size() << std::endl;

    return 0;
}

实际列队对于异步处理非常有帮助。假设有任务需要依次进入step1和step2两个阶段，如果想要提高任务的运行效率可以使用异步的方法，即step1产生结果后送入列队，然后开始新的step1。step2则会循环判断列队是否有值，若有则获取队列中的值执行step2。相当于在step1和step2中间建立了一个蓄水池，step1不断将结果扔向蓄水池，step2则在蓄水池有水（值）时执行step2。这样大大提高了整个任务的效率。

2、std::queue可以存储哪些数据类型？

std::queue 是 C++ 标准库中的容器之一，它可以存储任何 C++ 数据类型，包括内置数据类型（如 int、float、char 等）、自定义结构体、类对象等。当然除此之外我们还可以将指向数据的指针储存到queue中，不用关心指针指向的具体数据的类型、位置等。

现在我们要讨论一下，如果一个在step1产生的结果保存在cuda上，为了保证step1和step2的异步特性，如何将位于cuda上的数据放入列队？

2.1、queue如何存放位于cuda上的数据

我们需要知道，在使用显存时，通常需要先创建一个指针，然后分配显存并将其地址储存给指针：

// 假设在显存中会储存一个100个浮点数
int size=100
float *decodeed_cuda_buffer;

// 分配length * sizeof(float)个字节显存空间，并将空间地址储存在指针decodeed_cuda_buffer中
cudaMalloc((void **)(&decodeed_cuda_buffer), length * sizeof(float));

其中decodeed_cuda_buffer指针是位于host上的，而decodeed_cuda_buffer指向的内存是位于device上的。

所以我们可以在queue中储存指针decodeed_cuda_buffer，如图所示。首先通过cudaMalloc申请到新的显存后，将地址储存给指针*p；然后通过*p对显存数据进行操作（一般是写入新数据）；最后将指针*p送入queue。这样在queue中就储存了多个显存空间的地址，依次储存我们写入的数据。

实现代码如下，我们希望每次分配新的显存，并将地址给p_cuda_men；然后对cuda显存数据操作后，将其地址p_cuda_men送入queue。最终queue中存放多个cuda空间地址：

std::queue<float *> gpu_buffers_queue;  // 创建一个列队
float *p_cuda_men;  // 储存cuda显存地址
// 分配显存
for (int i=0;i<count;i++) {
    // part1：分配一块新的显存空间，并将地址储存在p_cuda_men中。
	cudaMalloc((void **)(&p_cuda_men), length * sizeof(float));
	operate(p_cuda_men)  // decodeed_cuda_buffer指向的显存的内容已改变
	// part2：将本次分配的显存块的地址放入列队
	gpu_buffers_queue.push(p_cuda_men);
}

一种错误的方法如下，显存只分配一次。可以看到for循环中只是修改了p_cuda_men指向的显存的值，p_cuda_men储存的地址没有改变。即循环push的是同一个显存地址，就是唯一一次cudaMalloc分配的显存地址。最终在gpu_buffers_queue中存放的是完全一样的显存地址，而显存中的值则是最后一次operate修改后的值：

std::queue<float *> gpu_buffers_queue;  // 创建一个列队
// 实例化一个float *并分配显存。decodeed_cuda_buffer指向分配内存的地址

// part1：创建新的指针，指向分配的显存。即该指针存放本次分配的显存块的地址。
float *p_cuda_men;
cudaMalloc((void **)(&p_cuda_men), length * sizeof(float));

// 分配显存
for (int i=0;i<count;i++) {
    // operate只是对唯一的一块显存进行操作
	operate(p_cuda_men)  
	// part2：将显存地送入列队
	gpu_buffers_queue.push(p_cuda_men);
}

2.2、如何从queue读取位于cuda上的数据？

看图，通过queue.front()可以获取到最先送入queue的指针p4，可以通过如cudaMemcpy(cpu_float_vector, p4, size * sizeof(float), cudaMemcpyDeviceToHost);等操作对该显存数据进行操作。

非常重要：cuda的显存分配后需要通过cudaFree(p4)手动释放。在queue中存放的是指向不同显存空间的指针，所以通过p=queue.front()获取到显存地址后，并对该地址数据进行相关操作，比如cudaFree§进行释放。否则只分配不释放，显存占用累计最终OOM。

如果是下面这种，最后没有cudaFree§。等下一次循环时，没有变量储存p4指向的空间，我们再也无法通过cudaFree()释放p4显存，导致一直占用:

loop:
	p=queue.front();  // 第一次循环：获取到*p4
	operate(p);  // 对p指向的显存数据进行某种需要的操作
	queue.pop();  // 弹出*p4，现在queue=[*p1, *p2, *p3]

合理的写法是：

loop:
	p=queue.front();  // 第一次循环：获取到*p4
	operate(p);  // 对p指向的显存数据进行某种需要的操作
	queue.pop();  // 弹出*p4，现在queue=[*p1, *p2, *p3]
	cudaFree(p);   // 释放p4指向的显存空间

2.3、注意：需要的最大显存

在上面的step1和step2异步中，step1向queue中写入，step2从queue读出。step1每写入依次代表分配了一块显存，而step2每读出一次则会销毁这块显存。

• 如果step1的处理速度小于step2，则step2经常性等待step1向queue中存入数据。所以分配的显存可能就只有1~2次（不严谨）。

• 如果step1的处理速度大于step2，则queue中存入比读取更快，导致queue的size越来越大，储存的指针数据越来越多。每个指针数据代表分配的一块显存，可想而知显存的占用也会越来越高（可用nvidia-smi查看）。所以实际运用中需要：1）多步任务中，需要合理分段，让各段的处理速度相当；2）限制queue长度，超过长度时新的数据直接覆盖queue的尾端（即最新放入的）。

在限制queue长度后，该程序需要使用的最大显存就是长度*每个显存块的size

3、一种更优的方法

上面的方法需要不断地创建释放显存，可能并不优雅。现在以下方法（理论上，未测试）优化：

• 创建一个cuda列队类，初始化的时候定义列队大小N并创建N个显存块。在使用时循环将N个显存块用来接收数据，即将旧数据覆盖储存新的数据即可。但是需要相应的逻辑代码控制列队的写和读是先入先出，且保证数据是有效的（防止读取的显存块是旧数据或者没有数据）。
• 使用共享内存。