OpenCV CUDA模块设备层-----逐通道最小值比较函数min()

• 操作系统：ubuntu22.04
• OpenCV版本：OpenCV4.9
• IDE:Visual Studio Code
• 编程语言：C++11

算法描述

OpenCV 的CUDA并行计算模块（cv::cudev）中的一个设备端内联函数，用于在CUDA核函数中对两个uchar2类型像素值进行逐通道最小值比较。

对两个 uchar2 类型变量（即 2 通道无符号字符类型）进行逐通道比较，返回每个通道中的最小值。
它常用于 GPU 加速的图像融合、像素比较、背景建模等场景。

函数原型

__device__ __forceinline__ uchar2 cv::cudev::min 	( 	const uchar2 &  	a,
		const uchar2 &  	b 
	) 	

参数

• a const uchar2& 第一个输入像素值（x, y 两个通道）
• b const uchar2& 第二个输入像素值

返回值

• 返回一个新的 uchar2 值；
• 每个通道的值为 a.x 和 b.x 的最小值，a.y 和 b.y 的最小值。

代码

#include <opencv2/cudaimgproc.hpp>
#include <opencv2/highgui.hpp>
#include <opencv2/cudev/util/vec_math.hpp>

using namespace cv;
using namespace cv::cudev;

// CUDA 核函数：取两个 uchar2 像素的最小值
template <typename Ptr1, typename Ptr2, typename DstPtr>
__global__ void minKernel(Ptr1 src1, Ptr2 src2, DstPtr dst, int width, int height) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    if (x < width && y < height) {
        // 获取两个图像的像素值
        uchar2 val1 = src1(y, x);
        uchar2 val2 = src2(y, x);

        // 取最小值并写入输出图像
        dst(y, x) = min(val1, val2);
    }
}

int main() {
    // 加载两张图像（假设尺寸相同），转换为 2 通道格式（比如合并两个灰度图为 uchar2）
    Mat h_img1 = imread("/media/dingxin/data/study/OpenCV/sources/images/img0.jpg", IMREAD_GRAYSCALE);
    Mat h_img2 = imread("/media/dingxin/data/study/OpenCV/sources/images/img1.jpg", IMREAD_GRAYSCALE);

    if (h_img1.empty() || h_img2.empty()) {
        std::cerr << "Failed to load images!" << std::endl;
        return -1;
    }

    // 转换为 2 通道图像（uchar2）
    Mat h_2ch[] = {h_img1, h_img2};
    Mat h_combined;
    merge(h_2ch, 2, h_combined);  // 合并为 CV_8UC2

    // 上传到 GPU
    cuda::GpuMat d_combined, d_result;
    d_combined.upload(h_combined);
    d_result.create(d_combined.size(), d_combined.type());

    // 构造访问器
    auto ptr = PtrStepSz<uchar2>(d_combined);
    auto dptr = PtrStepSz<uchar2>(d_result);

    // 设置核函数参数
    dim3 block(16, 16);
    dim3 grid((d_combined.cols + block.x - 1) / block.x,
              (d_combined.rows + block.y - 1) / block.y);

    // 调用核函数
    minKernel<<<grid, block>>>(ptr, ptr, dptr, d_combined.cols, d_combined.rows);
    cudaDeviceSynchronize();

    // 分离通道并显示结果
    Mat h_result;
    d_result.download(h_result);

    std::vector<Mat> channels;
    split(h_result, channels);

    imshow("Min Channel 1", channels[0]);
    imshow("Min Channel 2", channels[1]);
    waitKey(0);

    return 0;
}

运行结果