- 操作系統：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;
}
運行結果
