OpenCV CUDA模块设备层---- 绝对值函数abs()

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

算法描述

这是 OpenCV 的 cv::cudev 模块中用于 CUDA 设备端（device）的绝对值函数，专门处理 uchar1 类型（即单通道无符号字符）。

函数原型

__device__ __forceinline__ uchar1 cv::cudev::abs 	( 	const uchar1 &  	a	) 	

• device: 表示这个函数只能在 CUDA 设备端（GPU）运行。
• forceinline: 强制内联，提高性能。
• uchar1: CUDA 内建类型，表示一个包含 1 个 unsigned char 的向量（类似 struct { unsigned char x; };）。
• cv::cudev::abs(a): 返回输入值的绝对值。对于 uchar 来说其实等价于直接返回 a，因为它是无符号类型。

示例使用场景

在 CUDA 核函数中对图像像素进行操作时，可能会用到：

__global__ void absKernel(const cv::cudev::PtrStep<uchar> src,cv::cudev::PtrStep<uchar> dst)
{int x = blockIdx.x * blockDim.x + threadIdx.x;int y = blockIdx.y * blockDim.y + threadIdx.y;if (x < src.cols && y < src.rows){uchar val = src(y, x);dst(y, x) = cv::cudev::abs(make_uchar1(val)).x;}
}

示例代码

cu文件：

#include <opencv2/opencv.hpp>
#include <opencv2/cudaimgproc.hpp>
#include <opencv2/cudev/util/vec_math.hpp>using namespace cv;
using namespace cudev;__global__ void absKernel(PtrStepSz<uchar> src, PtrStepSz<uchar> dst)
{int x = blockIdx.x * blockDim.x + threadIdx.x;int y = blockIdx.y * blockDim.y + threadIdx.y;if (x < src.cols && y < src.rows){uchar1 val = make_uchar1(src(y, x));dst(y, x) = cv::cudev::abs(val).x;}
}int main()
{Mat h_src = imread("/media/dingxin/data/study/OpenCV/sources/images/Lenna.png", IMREAD_GRAYSCALE);if (h_src.empty()){std::cerr << "Failed to load image!" << std::endl;return -1;}cuda::GpuMat d_src, d_dst;d_src.upload(h_src);d_dst.create(h_src.size(), h_src.type());PtrStepSz<uchar> srcPtr(d_src);PtrStepSz<uchar> dstPtr(d_dst);dim3 block(16, 16);dim3 grid((srcPtr.cols + block.x - 1) / block.x,(srcPtr.rows + block.y - 1) / block.y);absKernel<<<grid, block>>>(srcPtr, dstPtr);cudaDeviceSynchronize();Mat h_dst;d_dst.download(h_dst);imshow("Original", h_src);imshow("Abs Result", h_dst);waitKey(0);return 0;
}

运行结果