CUDA实现图像的高斯滤波（opencv实现）

高斯滤波简介：
       高斯滤波是通过对输入数组的每个点与输入的高斯滤波模板执行卷积计算然后将这些结果一块组成了滤波后的输出数组，通俗的讲就是高斯滤波是对整幅图像进行加权平均的过程，每一个像素点的值都由其本身和邻域内的其他像素值经过加权平均后得到。

       高斯滤波的具体操作是：用一个模板（或称卷积、掩模）扫描图像中的每一个像素，用模板确定的邻域内像素的加权平均灰度值去替代模板中心像素点的值。

假设高斯函数的标准差δ=1，根据高斯函数可以得到二维函数

高斯模板生成

首先采用C++代码实现，结果如下：

#include <stdio.h>  
#include <iostream>
#include <windows.h>
#include <math.h>
using namespace std;
#define  PI      3.141592653
#define  sigma   1            //标准差值
#define  n       3        //卷积核（奇数）

int main(){

	double temp = 0;
	double h_sum = 0;  //归一化的总值  
	double distance = 0;

	double kernal[n][n];
	int center_x = n / 2, center_y = n / 2;
	cout << "未归一化的值：" << endl;
	for (int i = 0; i<n; i++)
	{
		for (int j = 0; j<n; j++)
		{
			distance = (i - center_x)*(i - center_x) + (j - center_x)*(j - center_x);
			temp = exp(-(distance) / (2 * sigma*sigma)) / (2 * PI*sigma*sigma);
			kernal[i][j] = temp;
			h_sum += temp;
			printf("%0.10f  ", temp);
		}
		cout << endl;
	}
	cout << endl;
	cout << "h_sum Value:" << h_sum << endl;
	cout << endl;
	cout << "归一化的值：" << endl;
	for (int i = 0; i<n; i++)
	{
		for (int j = 0; j<n; j++)
		{
			kernal[i][j] /= h_sum;
			printf("%0.10f  ", kernal[i][j]);
		}
		cout << endl;
	}
	system("pause");
	return 0;
}

运行结果：

CUDA核函数的实现：

__global__ void gaussianFilter(uchar *d_in, uchar *d_out, int width, int height)
{
	int tidx = blockDim.x * blockIdx.x + threadIdx.x;
	int tidy = blockDim.y * blockIdx.y + threadIdx.y;


	int sum = 0;
	int index = 0;


	if (tidx>2 && tidx<width - 2 && tidy>2 && tidy<height - 2)


	for (int m = tidx - 2; m < tidx + 3; m++)
	{
		for (int n = tidy - 2; n < tidy + 3; n++)
		{
			sum += d_in[m*width + n] * templates[index++];
		}
	}
	if (sum / 273<0)
		*(d_out + (tidx)*width + tidy) = 0;
	else if (sum / 273>255)
		*(d_out + (tidx)*width + tidy) = 255;
	else
		*(d_out + (tidx)*width + tidy) = sum / 273;


}

主函数编写步骤：

        读取待处理的图像；

        定义中间数据传递的指针并分配内存；

        将数据从Host端传到Device端；

        网格和块的分配；

        执行kernel函数；

        将数据从Device端传回到Host端；

        最后释放内存。

代码为：

int main()
{

	Mat srcImg = imread("D:\\pictures\\lena.jpg");
	Mat src; //   
	cvtColor(srcImg, src, CV_BGR2GRAY); // 转为灰度图像

	imshow("src_image", src);

	uchar *d_in;
	uchar *d_out;

	int width = srcImg.cols;
	int height = srcImg.rows;

	int memsize = width*height*sizeof(uchar);

	cudaMalloc((void**)&d_in, width * height * sizeof(uchar));
	cudaMalloc((void**)&d_out, width * height * sizeof(uchar));

	cudaMemcpy(d_in, src.data, memsize, cudaMemcpyHostToDevice);//数据从Host传到Device */

	int Gaussian[25] = { 1, 4, 7, 4, 1,
		4, 16, 26, 16, 4,
		7, 26, 41, 26, 7,
		4, 16, 26, 16, 4,
		1, 4, 7, 4, 1 };//总和为273  
	cudaMemcpyToSymbol(templates, Gaussian, 25 * sizeof(int));

	int bx = int(ceil((double)width / BLOCKDIM_X)); //网格和块的分配  
	int by = int(ceil((double)height / BLOCKDIM_Y));

	if (bx > GRIDDIM_X) bx = GRIDDIM_X;
	if (by > GRIDDIM_Y) by = GRIDDIM_Y;

	dim3 grid(bx, by);//网格的结构  
	dim3 block(BLOCKDIM_X, BLOCKDIM_Y);//块的结构  

	//kernel--高斯滤波  
	gaussianFilter << <grid, block >> >(d_in, d_out, width, height);

	cudaMemcpy(src.data, d_out, memsize, cudaMemcpyDeviceToHost);//数据传回主机端 */

	imshow("cuda_gaussian", src);
	imwrite("D:/gaussian.jpg", src);

	cudaFree(d_in);
	cudaFree(d_out);
	waitKey(0);
	return 0;
}

运行结果：