最近课程里面有用到 NVIDIA CUDA 框架进行并行编程,实现了一些非常基本的图像处理的操作。
使用 CUDA 实现的并行加速能够极大的提升图像处理的效率,这也是为什么近几年的深度学习框架都要依托于 CUDA 进行计算加速。CUDA 本质上是 C/C++ 的拓展,因此对 C/C++ 熟悉的话上手也会很快。
读取保存图像
C++ 中读取图像和保存图像不像 Python 那样方便简单。一开始我想用 OpenCV 来读取/写入图像,但对于我写一个小的示例程序来说可谓麻烦了些,涉及到编译的问题。
后来看到了 lodepng 这个小巧方便的库,支持 C 和 C++,使用时只需要引用一个头文件就可以。读取和保存文件只需要如下的几行代码:
1
2
3
4
5
6
7
8
9
10
11
#include"lodepng.h"// Read image
std::vector<unsignedchar>image;unsignedintwidth,height;unsignederror=lodepng::decode(image,width,height,input_file);// Save image
std::vector<unsignedchar>out_image(image.size(),255);error=lodepng::encode(output_file,out_image,width,height);
加载好的图片像素有 4 个通道,按照 RGBARGBA... 的顺序在一维数组中排列的,其中 A 表示的是透明度,0 表示全透明,255 则完全显示当前像素。
// CUDA Kernel
__global__voidRGB2GrayKernel(unsignedchar*input_image,unsignedchar*output_image,intwidth,intheight){// input_image size: width*height*Channels
// output_image size: width*height
introw=blockIdx.y*blockDim.y+threadIdx.y;intcol=blockIdx.x*blockDim.x+threadIdx.x;if((row<height)&&(col<width)){// Offset in Grayscale image
intoffset=row*width+col;// Get RGB values
unsignedcharr=input_image[offset*CHANNELS+0];unsignedcharg=input_image[offset*CHANNELS+1];unsignedcharb=input_image[offset*CHANNELS+2];// Convert to grayscale
output_image[offset]=0.2126f*r+0.7152f*g+0.0722f*b;}}
通过 CUDA Kernel 全局变量 blockDim、blockIdx 以及 threadIdx,可以得到当前运行 Kernel 的线程的编号,这样能够方便的访问部分数据。
在 CPU(host) 部分的 main 函数中,我们需要将图片数据拷贝到 device 上的 global memory 中,需要用到 cudaMalloc 以及 cudaMemcpy:
1
2
3
4
5
6
7
8
// Allocate memory for CUDA device
size_tmem_size=width*height*sizeof(unsignedchar);unsignedchar*dev_input,*dev_output;cudaMalloc(&dev_input,mem_size*CHANNELS);cudaMemcpy(dev_input,image.data(),mem_size*CHANNELS,cudaMemcpyHostToDevice);cudaMalloc(&dev_output,mem_size);
在 main 中调用 CUDA Kernel 的时候,需要用到尖括号的表达。下面在一个 Block 里我用了 16x16 个线程,然后对于 width x height 的图像,可以得到所需的 Grid 的 Size,使得整个图片能被所有的线程都处理到。
1
2
3
4
5
6
7
#define BLOCK_SIZE 16
// Invoke CUDA kernel
dim3dimBlock(BLOCK_SIZE,BLOCK_SIZE);dim3dimGrid(width/dimBlock.x,height/dimBlock.y);RGB2GrayKernel<<<dimGrid,dimBlock>>>(dev_input,dev_output,width,height);
在调用完 Kernel 之后,我们需要将 CUDA device 内存中的数据拷贝到 CPU 上 main 程序里,这样可以在后续使用 lodepng 将图片保存。最后不能忘记用 cudaFree 将分配的 GPU 显存给释放。
1
2
3
4
5
6
// Copy output from device
unsignedchar*image_y=newunsignedchar[width*height];cudaMemcpy(image_y,dev_output,mem_size,cudaMemcpyDeviceToHost);cudaFree(dev_input);cudaFree(dev_output);
#include"cuda_runtime.h"// Timing
cudaEvent_tstart,stop;cudaEventCreate(&start);cudaEventCreate(&stop);cudaEventRecord(start);//
// DO SOMETHING
// WITH OR WITHOUT CUDA
//
// Get elapsed time in ms
cudaEventRecord(stop);cudaEventSynchronize(stop);floatmilliseconds=0;cudaEventElapsedTime(&milliseconds,start,stop);
编译
编译需要用到 CUDA 库中的 NVIDIA CUDA Compiler (NVCC),在编译的时候注意 GPU 的架构,需要指定 -arch=compute_35 参数,具体的列表可以在 GPU Feature List 里查看。
下面我写了一个非常简单的 makefile 文件,写了一个 rgb2gray_cuda 的规则来使用系统中的 CUDA 10 下的 nvcc 来编译我的 rgb2gray.cu 源码。