- 操作系统：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;
}
运行结果
