CUDA-归一化算法

2024-10-22 08:47:14

实现原理

归一化算法的核心目的是将图像中的像素值调整到一个统一的范围，通常是[0, 255]。算法步骤如下：

找到图像中的最大值和最小值。在CUDA实现中，首先通过并行核函数getMaxMinValue_CUDA，利用atomicMax和atomicMin原子操作来确保线程安全地计算出最大和最小像素值。
在核函数 normalizeImage_CUDA中进行归一化处理。该过程通过将每个像素值减去最小值并除以最大值与最小值之差，最后乘以255，来线性映射像素值到新范围。

在图像处理中，归一化可以增强图像对比度，提高图像处理算法的稳定性和鲁棒性。通过将像素值调整到统一范围，可以更好地提取特征，促进后续的处理步骤，如边缘检测或分类。

本文将通过一个实战案例，进行归一化的展示。

（注意本文案例基于OpenCV实现，但是对CUDA而言，核心部分与OpenCV无关，可根据自身场景和数据结构进行更改。）

C++测试代码

ImageProcessing.cuh

#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;

#define TILE_WIDTH 16

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

// 图像归一化-CPU
cv::Mat normalizeImage_CPU(cv::Mat input);

// 图像归一化-GPU
cv::Mat normalizeImage_GPU(cv::Mat input);

ImageProcessing.cu

#include "ImageProcessing.cuh"

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

// 图像归一化-CPU
cv::Mat normalizeImage_CPU(cv::Mat input)
{
	int row = input.rows;
	int col = input.cols;
	// 确认最值
	uchar maxV = 0;
	uchar minV = 255;
	for (int i = 0; i < row; ++i)
	{
		for (int j = 0; j < col; ++j)
		{
			if (input.at<uchar>(i, j) > maxV)
			{
				maxV = input.at<uchar>(i, j);
			}
			if (input.at<uchar>(i, j) < minV)
			{
				minV = input.at<uchar>(i, j);
			}
		}
	}
	cout << "max value:" << int(maxV) << endl;
	cout << "min value:" << int(minV) << endl;
	// 归一化
	cv::Mat result = cv::Mat(row, col, CV_8UC1);
	for (int i = 0; i < row; ++i)
	{
		for (int j = 0; j < col; ++j)
		{
			result.at<uchar>(i, j) = uchar(float(input.at<uchar>(i, j) - minV) / (maxV - minV) * 255);
		}
	}
	return result;
}
// 获取最大最小值核函数
__global__ void getMaxMinValue_CUDA(uchar* inputImage, int width, int height, int *maxV, int *minV)
{
	int row = blockIdx.y * blockDim.y + threadIdx.y;
	int col = blockIdx.x * blockDim.x + threadIdx.x;

	if (row < height && col < width)
	{
		atomicMax(maxV, int(inputImage[row * width + col]));
		atomicMin(minV, int(inputImage[row * width + col]));
	}
}
// 图像归一化核函数
__global__ void normalizeImage_CUDA(uchar* inputImage, uchar* outputImage, int width, int height, uchar maxV, uchar minV)
{
	int row = blockIdx.y * blockDim.y + threadIdx.y;
	int col = blockIdx.x * blockDim.x + threadIdx.x;

	if (row < height && col < width)
	{
		outputImage[row * width + col] = uchar(float(inputImage[row * width + col] - minV) / (maxV - minV) * 255);
	}
}
// 图像归一化-GPU
cv::Mat normalizeImage_GPU(cv::Mat input)
{
	int row = input.rows;
	int col = input.cols;

	// 定义计时器
	float spendtime = 0.0f;
	cudaEvent_t start, end;
	cudaEventCreate(&start);
	cudaEventCreate(&end);

	// 分配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(TILE_WIDTH, TILE_WIDTH);
	dim3 gridSize((col + blockSize.x - 1) / blockSize.x, (row + blockSize.y - 1) / blockSize.y);

	// 求最值
	int h_maxValue = 0;
	int h_minValue = 255;
	int *d_maxValue;
	int *d_minValue;
	cudaMalloc((void**)&d_maxValue, sizeof(int));
	cudaMalloc((void**)&d_minValue, sizeof(int));
	cudaMemcpy(d_maxValue, &h_maxValue, sizeof(int), cudaMemcpyHostToDevice);
	cudaMemcpy(d_minValue, &h_minValue, sizeof(int), cudaMemcpyHostToDevice);
	getMaxMinValue_CUDA << <gridSize, blockSize >> > (d_inputImage, col, row, d_maxValue, d_minValue);
	cudaMemcpy(&h_maxValue, d_maxValue, sizeof(int), cudaMemcpyDeviceToHost);
	cudaMemcpy(&h_minValue, d_minValue, sizeof(int), cudaMemcpyDeviceToHost);
	cout << "max value:" << h_maxValue << endl;
	cout << "min value:" << h_minValue << endl;

	// 归一化
	cudaEventRecord(start, 0);
	normalizeImage_CUDA << <gridSize, blockSize >> > (d_inputImage, d_outputImage, col, row, h_maxValue, h_minValue);
	cudaEventRecord(end, 0);
	cudaEventSynchronize(end);
	cudaEventElapsedTime(&spendtime, start, end);
	cout << "kernel_CUDA spend time:" << spendtime << endl;

	// 将处理后的图像数据从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 "ImageProcessing.cuh"

void main()
{
    // 预准备
	warmupCUDA();

	cout << "normalizeImage test begin." << endl;
	// 加载
	cv::Mat src = imread("test pic/test5.jpg", 0);

	// 调整数据区间
	cv::Mat src2;
	cv::normalize(src, src2, 20, 230, NORM_MINMAX);

	// CPU版本
	clock_t s1, e1;
	s1 = clock();
	cv::Mat output1 = normalizeImage_CPU(src2);
	e1 = clock();
	cout << "CPU time:" << double(e1 - s1) << "ms" << endl;

	// GPU版本
	clock_t s2, e2;
	s2 = clock();
	cv::Mat output2 = normalizeImage_GPU(src2);
	e2 = clock();
	cout << "GPU time:" << double(e2 - s2) << "ms" << endl;

	// 查看输出
	cv::Mat test1 = output1.clone();
	cv::Mat test2 = output2.clone();

	cout << "normalizeImage test end." << endl;
	
}

测试效果

在本文案例中，我通过OpenCV自带的归一化函数将图像的最值设为20和230，如图1所示，图像对比度有所下降，然后用自写的归一化算法再将其调整后回0-255，图像对比度明显提升。速度方面，归一化核函数仅用了0.17ms，大多耗时还是在CPU和GPU的数据传输上。

如果函数有什么可以改进完善的地方，非常欢迎大家指出，一同进步何乐而不为呢~

如果文章帮助到你了，可以点个赞让我知道，我会很快乐~加油！

码农公寓

实现原理

C++测试代码

测试效果

相关文章