CUDA系列二、CUDA图像标准化

一、概述

在上一篇文章CUDA系列一、CUDA简介中我们对CUDA进行了介绍，了解到CUDA能让开发者利用GPU的强大算力进行大规模的并行计算，在本篇文章中将介绍如何使用CUDA进行图像标准化。

图像标准化在视觉领域内有广泛的应用，通过标准化能使图像数据有更稳定的分布，从而提高模型的训练效果和泛化能力，所以在图像相关的模型部署过程中，图像标准化通常是模型预处理的第一步，本篇通过图像标准化代码详细介绍并行带来的优势。

二、图像标准化原理

图像标准化，是将图像像素值调整到特定分布范围的操作，通常是把图像均值调整为0，标准差调整为1。对于一幅图像，其标准化公式为：

$$
I_{norm}(x,y) = \frac{I(x,y) - \mu}{\sigma}
$$

其中，$I(x,y)$ 是原始图像在坐标 $(x,y)$ 处的像素值，$\mu$ 是图像的均值，$\sigma$ 是图像的标准差。通俗地讲就是每个像素需要减去均值再除以标准差。

三、代码及解析

3.1 头文件

下面代码是 cu_standardization.hpp 头文件，使用OpenCV读取和存储图像，所以需要加载OpenCV相关的库，然后声明了一个 standardizationRGB 函数用于图像的标准化调用。

#pragma once
#include <array>
#include <memory>
#include <string>
#include <vector>
#include <opencv2/highgui/highgui.hpp>
#include <chrono>
#include <opencv2/opencv.hpp>
#include <device_launch_parameters.h>
#include "opencv2/core/cuda/common.hpp"

// 标准化
void standardizationRGB(const cv::cuda::GpuMat cu_src, cv::cuda::GpuMat cu_dst);

3.2 标准化核心文件

下面代码是 cu_standardization.cu 文件，首先通过 __constant__ 常量内存定义了RGB图像的均值和方差，常量内存在核函数执行期间不会发生变化，且数据在所有线程共享。然后在 __global__ void kenlStandardizationRGB 核函数中，根据均值和标准差对每个像素进行标准化处理。由于是RGB图像，所以每个线程独立负责某一坐标下RGB三个像素值的计算。

void standardizationRGB 函数定义了线程块的大小为 BLOCK_SIZE x BLOCK_SIZE = 1024，cv::cuda::device::divUp 用于向上取整的除法运算，所以开辟的总线程数量不小于图像的宽乘高。cv::cuda::PtrStepSz 是OpenCV中用于在CUDA设备上处理二维数组的类，它封装了二维数组的行数、列数、数据指针以及每行的步长信息，我们将图像传入到这个类中便于核函数的调用。

// cu文件
#include "cu_standardization.hpp"

// DEBUG 不支持过大的BLOCK
#ifdef _DEBUG
#define BLOCK_SIZE 16
#else
#define BLOCK_SIZE 32
#endif

// 定义一个极小值，避免除零
#define EPSILON 1e-8

// 定义常量内存均值和方差
__constant__ float3 mean_ = { 0.485f, 0.456f, 0.406f };
__constant__ float3 std_ = { 0.229f, 0.224f, 0.225f };

// 计算
__global__ void kenlStandardizationRGB(
    const cv::cuda::PtrStepSz<uchar3> src,
    cv::cuda::PtrStepSz<float3> dst)
{
    const int x = blockIdx.x * blockDim.x + threadIdx.x;
    const int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x >= src.cols || y >= src.rows)
        return;

    dst(y, x).x = (src(y, x).x / 255.0 - mean_.x) / (std_.x + EPSILON);
    dst(y, x).y = (src(y, x).y / 255.0 - mean_.y) / (std_.y + EPSILON);
    dst(y, x).z = (src(y, x).z / 255.0 - mean_.z) / (std_.z + EPSILON);
}

void standardizationRGB(const cv::cuda::GpuMat cu_src, cv::cuda::GpuMat cu_dst)
{
    dim3 block(BLOCK_SIZE, BLOCK_SIZE);
    dim3 grid(
        cv::cuda::device::divUp(cu_src.cols, block.x),
        cv::cuda::device::divUp(cu_src.rows, block.y));

    cv::cuda::PtrStepSz<float3> dst_(
        cu_dst.rows, cu_dst.cols, (float3 *) cu_dst.data, cu_dst.step);
    cv::cuda::PtrStepSz<uchar3> src_(
        cu_src.rows, cu_src.cols, (uchar3 *) cu_src.data, cu_src.step);

    // 标准化
    kenlStandardizationRGB<<<grid, block>>>(src_, dst_);
    cudaSafeCall(cudaDeviceSynchronize());
}

3.3 主函数

下面代码是执行的主函数 main.cpp。这个文件加载 cu_standardization.hpp 头文件，并使用OpenCV读取一张彩色图像，再使用 cv::cuda::GpuMat 将主机的图像传入到设备，也就是从CPU中的Mat传到GPU中的GpuMat。然后定义一个和输入图像一样大小的GPU图像作为输出值，调用 standardizationRGB 对图像进行标准化，最后将标准化的图像从GPU中拷贝到CPU中。

#include <iostream>
#include "cu_standardization.hpp"

int main()
{
    std::string path = R"(D:\img.png)";
    cv::Mat img_src = cv::imread(path);
    cv::cuda::GpuMat gpu_img_src;
    gpu_img_src.upload(img_src);

    cv::cuda::GpuMat gpu_img_dst(img_src.size(), CV_32FC3);
    standardizationRGB(gpu_img_src, gpu_img_dst);

    cv::Mat img_dst;
    gpu_img_dst.download(img_dst);
}

这就是整个GPU处理图像标准化的全过程，希望这篇文章能带你正式走进CUDA的开发之路。

本文仅供学习参考，未经允许，拒绝转载或者用作他用。