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

• 操作系统：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,               // 注意：不是 const
    PtrStep<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;

    // 上传到 GPU
    cuda::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);

    // 启动 kernel
    copyWarpCopyKernel<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;
}

运行结果