CUDA学习笔记（对图像进行并行加速运算）

在对CUDA的kernel函数有了一定了解之后，就可以对图像进行多线程加速计算，下面是我在CPU和GPU上各自遍历一张图像，对比两者的性能差异。
这里需要用到opencv的函数读取图像以及显示图像，因此opencv的头文件需要包含进去。其他的注意事项主要是图像的内存分配和存储类型问题，比如彩色图像是三通道的，申请的内存为row * col * sizeof(uchar3)。

首先在CPU上遍历一张彩色图像的所有像素：

	for (int i = 0; i < row; i++)
	{
		for (int j = 0; j < col; j++)
		{
			int index = i * row + j;			//像素索引
		    dst.data[3 * index + 0] = (uchar)src.data[3 * index + 0];
			dst.data[3 * index + 1] = (uchar)src.data[3 * index + 1];
			dst.data[3 * index + 2] = (uchar)src.data[3 * index + 2];

		}
	}

在GPU上遍历一张图像的所有像素，其kernel函数为：

__global__ void Traverse(uchar3* _src_dev, uchar3* dst_dev, int col,int row )
{
	// 一维数据索引计算（万能计算方法）
	int tid = blockIdx.z * (gridDim.x * gridDim.y) * (blockDim.x * blockDim.y * blockDim.z) \
		+ blockIdx.y * gridDim.x * (blockDim.x * blockDim.y * blockDim.z) \
		+ blockIdx.x * (blockDim.x * blockDim.y * blockDim.z) \
		+ threadIdx.z * (blockDim.x * blockDim.y) \
		+ threadIdx.y * blockDim.x \
		+ threadIdx.x;
	if (tid < col * row* 3) {
		dst_dev[tid].x = _src_dev[tid].x ;
		dst_dev[tid].y = _src_dev[tid].y ;
		dst_dev[tid].z = _src_dev[tid].z ;
	}
}

线程索引计算方法各有不同，很多新手会在这个地方犯迷糊，从其他地方看到了这个万能计算方法（仅限于单个维度的线程索引），用起来相当省事。
下面是完整的代码：

#include "cuda_runtime.h"
#include "iostream"
#include "opencv2/opencv.hpp"
#include <windows.h>   

using namespace std;
using namespace cv;

__global__ void Traverse(uchar3* _src_dev, uchar3* dst_dev, int col,int row )
{
		//int tid = threadIdx.x;		
		//for (int i = tid; i < col*row; i += 32) {
		//	dst_dev[i] = _src_dev[i];
		//}

	// 一维数据索引计算（万能计算方法）
	int tid = blockIdx.z * (gridDim.x * gridDim.y) * (blockDim.x * blockDim.y * blockDim.z) \
		+ blockIdx.y * gridDim.x * (blockDim.x * blockDim.y * blockDim.z) \
		+ blockIdx.x * (blockDim.x * blockDim.y * blockDim.z) \
		+ threadIdx.z * (blockDim.x * blockDim.y) \
		+ threadIdx.y * blockDim.x \
		+ threadIdx.x;
	if (tid < col * row* 3) {
		dst_dev[tid].x = _src_dev[tid].x ;
		dst_dev[tid].y = _src_dev[tid].y ;
		dst_dev[tid].z = _src_dev[tid].z ;
	}
}

int main()
{
	Mat src = cv::imread("E:\\picture\\source\\lena.jpg");
	Mat dst=src;
	int col = src.cols;
	int row=src.rows;
	//cv::imshow("src", src);
	//CPU遍历像素值
	clock_t start = clock();
	for (int i = 0; i < row; i++)
	{
		for (int j = 0; j < col; j++)
		{
			int index = i * row + j;			//像素索引
		    dst.data[3 * index + 0] = (uchar)src.data[3 * index + 0];
			dst.data[3 * index + 1] = (uchar)src.data[3 * index + 1];
			dst.data[3 * index + 2] = (uchar)src.data[3 * index + 2];

		}
	}
	cout << "cpu所耗费的时间：" << (double)(clock() - start)<< "ms" << "\n";
	cv::imshow("dst", dst);

	//GPU遍历像素值
	
	uchar3* src_dev, * dst_dev;
	Mat _dst = Mat(row, col , CV_8UC3);
	cudaMalloc((void**)&src_dev, row * col * sizeof(uchar3));
	cudaMalloc((void**)&dst_dev, row * col * sizeof(uchar3));
	cudaMemcpy(src_dev, src.data, row * col * sizeof(uchar3), cudaMemcpyHostToDevice);

	dim3 grid(1 + (col * row / (32 * 32 + 1)), 1, 1);      // grid
	dim3 block(32, 32, 1);
	//clock_t start2 = clock();
	cudaEvent_t start1, stop1;
	cudaEventCreate(&start1);
	cudaEventCreate(&stop1);
	cudaEventRecord(start1, 0);

	Traverse << < grid, block >> > (src_dev, dst_dev,col,row);
	cudaEventRecord(stop1, 0);
	cudaEventSynchronize(stop1);
	float time1;
	cudaEventElapsedTime(&time1, start1, stop1);
	cout << "Gpu所耗费的时间：" << time1 << "ms" << "\n";
	cudaMemcpy(_dst.data, dst_dev, row * col * sizeof(uchar3), cudaMemcpyDeviceToHost);
	cv::imshow("_dst", _dst);

	// free
	cudaFree(src_dev);
	cudaFree(dst_dev);

	waitKey(0);
	system("pause");
	return 0;
}

最后的计算结果为：

需要注意的是，由于一张图片的计算量有限，除非是特别大的图像或者数据集，否则cuda并行计算的优势并不明显，甚至会比CPU的计算时间还要长，此外这里是对像素进行遍历拷贝复制，需要频繁的进行内存读取操作，这并不能体现出GPU的并行计算优势。
刚才演示的是单个维度的线程索引，现在用两个维度的线程进行测试，核函数改为：

__global__ void Traverse(uchar3* _src_dev, uchar3* dst_dev, int col, int row)
{
	int x = threadIdx.x + blockIdx.x * blockDim.x;
	int y = threadIdx.y + blockIdx.y * blockDim.y;
	int offset = x + y * blockDim.x * gridDim.x;

	dst_dev[offset].x = _src_dev[offset].x;
	dst_dev[offset].y = _src_dev[offset].x;
	dst_dev[offset].z = _src_dev[offset].x;
}

上面的测试用例是512*512的lena图像，使用CUDA的时间优势还是没有体现出来，为了证明GPU的加速性能，下面换用3000*2000的图片进行测试。

这的例子只是为了介绍一下最基本的CUDA对图片进行并行加速的思路，方便大家举一反三，触类旁通，其实在这个基础上稍加改动就可以尝试用CUDA对图像进行各种图像处理的操作，比如阈值分割，两个图象叠加等。
比如核函数改为：

__global__ void Traverse(uchar3* _src_dev, uchar3* dst_dev, int col,int row )
{
	// 一维数据索引计算（万能计算方法）
	int tid = blockIdx.z * (gridDim.x * gridDim.y) * (blockDim.x * blockDim.y * blockDim.z) \
		+ blockIdx.y * gridDim.x * (blockDim.x * blockDim.y * blockDim.z) \
		+ blockIdx.x * (blockDim.x * blockDim.y * blockDim.z) \
		+ threadIdx.z * (blockDim.x * blockDim.y) \
		+ threadIdx.y * blockDim.x \
		+ threadIdx.x;
	if (tid < col * row * 3) {
		dst_dev[tid].x = 255- _src_dev[tid].x;
		dst_dev[tid].y = 255- _src_dev[tid].y;
		dst_dev[tid].z = 255- _src_dev[tid].z;
	}
}

结果为负片效果：

核函数改为：

__global__ void Traverse(uchar3* _src_dev, uchar3* dst_dev, int col,int row )
{
	// 一维数据索引计算（万能计算方法）
	int tid = blockIdx.z * (gridDim.x * gridDim.y) * (blockDim.x * blockDim.y * blockDim.z) \
		+ blockIdx.y * gridDim.x * (blockDim.x * blockDim.y * blockDim.z) \
		+ blockIdx.x * (blockDim.x * blockDim.y * blockDim.z) \
		+ threadIdx.z * (blockDim.x * blockDim.y) \
		+ threadIdx.y * blockDim.x \
		+ threadIdx.x;
	if (tid < col * row * 3) {
		dst_dev[tid].x = 0.3 * _src_dev[tid].x;
		dst_dev[tid].y = 0.6 * _src_dev[tid].y;
		dst_dev[tid].z = 0.1 * _src_dev[tid].z;
	}
}

更改后效果如下 ：

核函数改为：

__global__ void Traverse(uchar3* _src_dev, uchar* dst_dev, int col,int row )
{
	// 一维数据索引计算（万能计算方法）
	int tid = blockIdx.z * (gridDim.x * gridDim.y) * (blockDim.x * blockDim.y * blockDim.z) \
		+ blockIdx.y * gridDim.x * (blockDim.x * blockDim.y * blockDim.z) \
		+ blockIdx.x * (blockDim.x * blockDim.y * blockDim.z) \
		+ threadIdx.z * (blockDim.x * blockDim.y) \
		+ threadIdx.y * blockDim.x \
		+ threadIdx.x;
	if (tid < col * row * 3) {
		dst_dev[tid] = 0.3 * _src_dev[tid].x+ 0.6 * _src_dev[tid].y+ 0.1 * _src_dev[tid].z;
	}

}

更改后效果如下 ：