cuda图像处理_CUDA随笔之Stream的使用

出于个人兴趣和工作需要，最近接触了GPU编程。于是想写篇文章（或一系列文章），总结一下学习所得，防止自己以后忘了。

这篇短文主要介绍CUDA里面Stream的概念。用到CUDA的程序一般需要处理海量的数据，内存带宽经常会成为主要的瓶颈。在Stream的帮助下，CUDA程序可以有效地将内存读取和数值运算并行，从而提升数据的吞吐量。

本文使用了一个非常naive的图像处理例子：像素色彩空间转换，将一张7680x4320的8-bit BRGA图像转成同样尺寸的8-bit YUV。计算非常简单，就是数据量非常大。转换公式直接照抄维基百科（https://en.wikipedia.org/wiki/YUV#Conversion_to/from_RGB）

由于GPU和CPU不能直接读取对方的内存，CUDA程序一般会有一下三个步骤：1）将数据从CPU内存转移到GPU内存，2）GPU进行运算并将结果保存在GPU内存，3）将结果从GPU内存拷贝到CPU内存。

如果不做特别处理，那么CUDA会默认只使用一个Stream（Default Stream）。在这种情况下，刚刚提到的三个步骤就如菊花链般蛋疼地串联，必须等一步完成了才能进行下一步。是不是很别扭？（短文末尾附有完整代码）

数值计算必须等数据拷贝完全结束后才开始能

uint8_t* bgraBuffer;
uint8_t* yuvBuffer;
uint8_t* deviceBgraBuffer;
uint8_t* deviceYuvBuffer;

const int dataSizeBgra = 7680 * 4320 * 4;
const int dataSizeYuv = 7680 * 4320 * 3;

cudaMallocHost(&bgraBuffer, dataSizeBgra);
cudaMallocHost(&yuvBuffer, dataSizeYuv);
cudaMalloc(&deviceBgraBuffer, dataSizeBgra);
cudaMalloc(&deviceYuvBuffer, dataSizeYuv);

//随机生成8K的BGRA图像
GenerateBgra8K(bgraBuffer, dataSizeBgra);

//将图像拷贝到GPU内存
cudaMemcpy(deviceBgraBuffer, bgraBuffer, dataSizeBgra, cudaMemcpyHostToDevice);

//CUDA kernel将 BGRA 转换为 YUV
convertPixelFormat<<<4096, 1024>>>(deviceBgraBuffer, deviceYuvBuffer, 7680*4320);

//等待数值计算完成
cudaDeviceSynchronize()

//将转换完的图像拷贝回CPU内存
cudaMemcpy(yuvBuffer, deviceYuvBuffer, dataSizeYuv, cudaMemcpyDeviceToHost);

cudaFreeHost(bgraBuffer);
cudaFreeHost(yuvBuffer);
cudaFree(deviceBgraBuffer);
cudaFree(deviceYuvBuffer);

NVIDIA家的GPU有一下很不错的技能（不知道是不是独有）：

1. 数据拷贝和数值计算可以同时进行。
2. 两个方向的拷贝可以同时进行（GPU到CPU，和CPU到GPU），数据如同行驶在双向快车道。

但同时，这数据和计算的并行也有一点合乎逻辑的限制：进行数值计算的kernel不能读写正在被拷贝的数据。

Stream正是帮助我们实现以上两个并行的重要工具。基本的概念是：

1. 将数据拆分称许多块，每一块交给一个Stream来处理。
2. 每一个Stream包含了三个步骤：1）将属于该Stream的数据从CPU内存转移到GPU内存，2）GPU进行运算并将结果保存在GPU内存，3）将该Stream的结果从GPU内存拷贝到CPU内存。
3. 所有的Stream被同时启动，由GPU的scheduler决定如何并行。

在这样的骚操作下，假设我们把数据分成A，B两块，各由一个Stream来处理。A的数值计算可以和B的数据传输同时进行，而A与B的数据传输也可以同时进行。由于第一个Stream只用到了数据A，而第二个Stream只用到了数据B，“进行数值计算的kernel不能读写正在被拷贝的数据”这一限制并没有被违反。效果如下：

用2个Stream（上）与只用一个Default Stream（下）的对比

实际上在NSight Profiler里面看上去是这样（这里用了8个Stream）：