CUDA-均值滤波算法

作者：翟天保Steven
版权声明：著作权归作者所有，商业转载请联系作者获得授权，非商业转载请注明出处

实现原理

均值滤波是一种常见的图像处理方法，用于去除图像中的噪声。其原理很简单：对图像中的每个像素，取其周围邻域像素的均值作为新的像素值。这个邻域可以是一个矩形、圆形或其他形状的区域，取决于具体的实现。

具体来说，均值滤波的原理可以分为以下几个步骤：

定义滤波器大小：确定用于计算均值的邻域大小。通常使用一个固定大小的窗口，例如3x3、5x5或7x7的矩形窗口。
遍历图像：对于图像中的每个像素，将邻域范围内的所有像素值相加。
计算均值：将邻域内所有像素值的总和除以像素数量，得到均值。
更新像素值：用计算得到的均值替换原始像素的值。

这个过程会对图像中的每个像素都进行处理，从而实现对整个图像的平滑处理，去除噪声。

均值滤波的优点是简单易实现，并且能有效地去除图像中的噪声。然而，它也有一些缺点，比如可能会导致图像边缘模糊，因为它只是简单地取邻域内像素的平均值，而不考虑像素之间的空间关系。

本文主要目的在于展示CUDA版本的性能提升效果，采用常规思路实现，CPU版本应用了并行提速，与GPU并行客观对比。

功能函数代码

// 均值滤波核函数
__global__ void meanFilter_CUDA(uchar* inputImage, uchar* outputImage, int width, int height, int windowSize)
{
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;

    if (row < height && col < width) 
    {
        int r = windowSize / 2;
        int ms = max(row - r, 0);
        int me = min(row + r, height - 1);
        int ns = max(col - r, 0);
        int ne = min(col + r, width - 1);

        float sum = 0.0f;
        int count = 0;

        for (int m = ms; m <= me; ++m) 
        {
            for (int n = ns; n <= ne; ++n) 
            {
                sum += inputImage[m * width + n];
                count++;
            }
        }
        outputImage[row * width + col] = uchar(sum / count);
    }
}

C++测试代码

Filter.h

#pragma once
#include <cuda_runtime.h>
#include <iostream>
#include <vector>
#include <opencv2/opencv.hpp>
#include <device_launch_parameters.h>

using namespace cv;
using namespace std;

// 预准备过程
void warmupCUDA();

// 均值滤波-GPU
cv::Mat filterAverage_GPU(cv::Mat input, int FilterWindowSize);

Filter.cu

#include "Filter.h"

// 预准备过程
void warmupCUDA()
{
    float* dummy_data;
    cudaMalloc((void**)&dummy_data, sizeof(float));
    cudaFree(dummy_data);
}

// 均值滤波核函数
__global__ void meanFilter_CUDA(uchar* inputImage, uchar* outputImage, int width, int height, int windowSize)
{
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;

    if (row < height && col < width) 
    {
        int r = windowSize / 2;
        int ms = max(row - r, 0);
        int me = min(row + r, height - 1);
        int ns = max(col - r, 0);
        int ne = min(col + r, width - 1);

        float sum = 0.0f;
        int count = 0;

        for (int m = ms; m <= me; ++m) 
        {
            for (int n = ns; n <= ne; ++n) 
            {
                sum += inputImage[m * width + n];
                count++;
            }
        }
        outputImage[row * width + col] = uchar(sum / count);
    }
}

// 均值滤波-GPU
cv::Mat filterAverage_GPU(cv::Mat input, int FilterWindowSize)
{
    int row = input.rows;
    int col = input.cols;

    // 分配GPU内存
    uchar* d_inputImage, * d_outputImage;
    cudaMalloc(&d_inputImage, row * col * sizeof(uchar));
    cudaMalloc(&d_outputImage, row * col * sizeof(uchar));

    // 将输入图像数据从主机内存复制到GPU内存
    cudaMemcpy(d_inputImage, input.data, row * col * sizeof(uchar), cudaMemcpyHostToDevice);

    // 计算块和线程的大小
    dim3 blockSize(16, 16);
    dim3 gridSize((col + blockSize.x - 1) / blockSize.x, (row + blockSize.y - 1) / blockSize.y);

    // 调用CUDA内核
    meanFilter_CUDA << <gridSize, blockSize >> > (d_inputImage, d_outputImage, col, row, FilterWindowSize);

    // 将处理后的图像数据从GPU内存复制回主机内存
    cv::Mat output(row, col, CV_8UC1);
    cudaMemcpy(output.data, d_outputImage, row * col * sizeof(uchar), cudaMemcpyDeviceToHost);

    // 清理GPU内存
    cudaFree(d_inputImage);
    cudaFree(d_outputImage);

    return output;
}

main.cpp

#include "Filter.h"

// 均值滤波-CPU
cv::Mat filterAverage_CPU(cv::Mat input, int FilterWindowSize)
{
	int row = input.rows;
	int col = input.cols;

	// 预设输出
	cv::Mat output = input.clone();

	// 均值滤波
	int r = FilterWindowSize / 2;
#pragma omp parallel for
	for (int i = 0; i < row; ++i)
	{
		for (int j = 0; j < col; ++j)
		{
			// 卷积窗口边界限制，防止越界
			int ms = ((i - r) > 0) ? (i - r) : 0;
			int me = ((i + r) < (row - 1)) ? (i + r) : (row - 1);
			int ns = ((j - r) > 0) ? (j - r) : 0;
			int ne = ((j + r) < (col - 1)) ? (j + r) : (col - 1);
			// 求窗口内有效数据的均值
			int count = 0;
			float sum = 0.0f;
			for (int m = ms; m <= me; ++m)
			{
				for (int n = ns; n <= ne; ++n)
				{
					sum += input.at<uchar>(m, n);
					count++;
				}
			}
			output.at<uchar>(i, j) = uchar(sum / count);
		}
	}

	return output;
}

void main()
{
	// 加载
	cv::Mat src = imread("test1.jpg", 0);
	int winSize = 51;
	// 预准备过程
	warmupCUDA();
	cout << "filterWindowSize:" << winSize << endl;
	cout << "size: " << src.cols << " * " << src.rows << endl;

	// CPU版本
	clock_t s1, e1;
	s1 = clock();
	cv::Mat output1 = filterAverage_CPU(src, winSize);
	e1 = clock();
	cout << "CPU time:" << double(e1 - s1) / 1000 << endl;

	// GPU版本
	clock_t s2, e2;
	s2 = clock();
	cv::Mat output2 = filterAverage_GPU(src, winSize);
	e2 = clock();
	cout << "GPU time:" << double(e2 - s2) / 1000 << endl;

	// 检查
	int row = src.rows;
	int col = src.cols;
	bool flag = true;
	for (int i = 0; i < row; ++i)
	{
		for (int j = 0; j < col; ++j)
		{
			if (output1.at<uchar>(i, j) != output2.at<uchar>(i, j))
			{
				flag = false;
				break;
			}
		}
		if (!flag)
		{
			break;
		}
	}
	if (flag)
	{
		cout << "ok!" << endl;
	}
	else
	{
		cout << "error!" << endl;
	}

	// 查看输出
	cv::Mat test1 = output1.clone();
	cv::Mat test2 = output2.clone();
	cout << "mean test." << endl;
	
}