文章目录
- 1、std::queue列队如何实现异步？
- 2、std::queue可以存储哪些数据类型？
- 2.1、queue如何存放位于cuda上的数据
- 2.2、如何从queue读取位于cuda上的数据？
- 2.3、注意：需要的最大显存
- 3、一种更优的方法
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个显存块用来接收数据，即将旧数据覆盖储存新的数据即可。但是需要相应的逻辑代码控制列队的写和读是先入先出，且保证数据是有效的（防止读取的显存块是旧数据或者没有数据）。
- 使用共享内存。
