当前位置：首页 > news >正文

OpenCV CUDA模块设备层-----在 GPU 上执行类似于 std::copy 的操作函数warpCopy()

news 2025/7/4 6:55:41

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

算法描述

OpenCV 的 CUDA 模块（cudev）中的一个设备端内联模板函数，用于在 GPU 上执行类似于 std::copy 的操作，但专门针对 warp 规模的数据复制。
该函数的作用是：
将一个范围内的元素从输入迭代器 beg 到 end 之间复制到输出迭代器 out 所指向的位置。

函数原型

 template<class InIt , class OutIt >
__device__ __forceinline__ OutIt cv::cudev::warpCopy 	
( 	InIt  	beg,InIt  	end,OutIt  	out 
)

参数

InIt 输入迭代器类型（例如 PtrTraits<…>::ptr_type）
OutIt 输出迭代器类型（例如 PtrTraits<…>::ptr_type）

返回值

返回最终的输出迭代器 out，指向最后一个复制元素之后的位置，便于链式调用或后续操作。

使用场景

这个函数通常用于以下情况：

在 CUDA kernel 中进行快速内存拷贝（如图像像素、数组等）
实现自定义的图像变换或数据搬运逻辑
构建更复杂的并行算法（如分块处理、扫描、归约）

它非常适合在每个线程负责多个数据项的场景下使用（即“warp-level”粒度的复制），可以提高内存访问效率和并行利用率。

代码

#include <opencv2/opencv.hpp>
#include <opencv2/cudaimgproc.hpp>
#include <opencv2/cudev.hpp>using namespace cv;
using namespace cv::cudev;// 使用 warpCopy 的 kernel，用于高效复制一行像素
template <typename T>
__global__ void copyWarpCopyKernel(PtrStep<T> src,               // 注意：不是 constPtrStep<T> dst,int roiX, int roiY,int roiWidth, int roiHeight)
{int y = blockIdx.y * blockDim.y + threadIdx.y;if (y < roiHeight) {T* srcRow = &src(roiY + y, roiX);  // 正确获取非 const 指针T* dstRow = &dst(y, 0);warpCopy(srcRow, srcRow + roiWidth, dstRow);}
}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;}// 设置 ROI 参数int roiX = 100;int roiY = 50;int roiWidth = 320;int roiHeight = 240;// 上传到 GPUcuda::GpuMat d_src, d_dst;d_src.upload(h_src);d_dst.create(roiHeight, roiWidth, d_src.type());// 配置 kernel 参数（仅在 Y 方向并行）dim3 block(16, 16);dim3 grid(1, (roiHeight + block.y - 1) / block.y);// 启动 kernelcopyWarpCopyKernel<uchar><<<grid, block>>>(d_src, d_dst, roiX, roiY, roiWidth, roiHeight);// 下载结果Mat h_dst;d_dst.download(h_dst);// 显示结果imshow("Copied ROI", h_dst);waitKey(0);return 0;
}