- 操作系统：ubuntu22.04
- OpenCV版本：OpenCV4.9
- IDE:Visual Studio Code
- 编程语言：C++11
算法描述
在同一个线程块（thread block内，将 [beg, end) 范围内的数据并行地复制到 out 开始的位置。
它使用了 CUDA 线程协作机制（warp-level 或 block-level） 来实现高效的块级拷贝，通常比简单的逐线程拷贝更快。
函数原型
_device__ static __forceinline__ void cv::cudev::blockCopy
(InIt beg,InIt end,OutIt out
)
参数
| 参数名 | 类型 | 含义 |
|---|---|---|
| beg | InIt | 输入范围的起始迭代器（或指针） |
| end | InIt | 输入范围的结束迭代器（不包含该位置） |
| out | OutIt | 输出范围的起始迭代器（或指针） |
使用场景
- 图像处理中的 局部块拷贝
- 构建 GPU 并行归约（reduction）算法前的数据准备
- 实现滑动窗口（sliding window）或卷积操作时的 tile 数据加载
- 需要多个线程协同完成连续内存拷贝的任务
示例代码
#include <opencv2/cudev/functional/functional.hpp>
#include <cstdio>
#include <opencv2/core.hpp>
#include <opencv2/core/cuda.hpp>
#include <opencv2/cudev/block/block.hpp>
#include <iostream>using namespace cv::cudev;// 定义 tile 大小为 16x16
#define TILE_SIZE 16__global__ void processImageKernel(const uchar* input, size_t pitchIn,uchar* output, size_t pitchOut,int width, int height) {// 共享内存用于存储当前线程块处理的 tile__shared__ uchar tile[TILE_SIZE][TILE_SIZE];// 当前线程负责的图像坐标int x = blockIdx.x * TILE_SIZE + threadIdx.x;int y = blockIdx.y * TILE_SIZE + threadIdx.y;// 检查是否在图像范围内if (x >= width || y >= height)return;// 输入和输出指针const uchar* in_ptr = input + y * pitchIn + x;uchar* out_ptr = output + y * pitchOut + x;// 将当前 tile 拷贝到共享内存blockCopy(in_ptr, in_ptr + TILE_SIZE * TILE_SIZE, &tile[0][0]);__syncthreads(); // 必须同步所有线程，确保共享内存拷贝完成// 对 tile 中的数据进行简单处理（比如增加亮度）uchar value = tile[threadIdx.y][threadIdx.x];value = min(value + 50, (uchar)255);// 写回到输出位置blockCopy(&value, &value + 1, out_ptr);
}int main() {// 创建一个测试图像（8UC1）cv::Mat h_src = cv::Mat::zeros(128, 128, CV_8UC1);h_src.setTo(100); // 填充为灰度值 100// 上传到 GPUcv::cuda::GpuMat d_src, d_dst;d_src.upload(h_src);d_dst.create(h_src.size(), h_src.type());// 设置 kernel 参数dim3 threads(TILE_SIZE, TILE_SIZE);dim3 blocks((h_src.cols + TILE_SIZE - 1) / TILE_SIZE,(h_src.rows + TILE_SIZE - 1) / TILE_SIZE);// 调用 kernelprocessImageKernel<<<blocks, threads>>>(d_src.ptr<uchar>(), d_src.step,d_dst.ptr<uchar>(), d_dst.step,d_src.cols, d_dst.rows);// 下载结果cv::Mat h_dst;d_dst.download(h_dst);// 显示部分结果std::cout << "Output image sample:\n" << h_dst(cv::Rect(0, 0, 5, 5)) << std::endl;return 0;
}
运行结果
Output image sample:
[150, 0, 0, 0, 0;0, 0, 0, 0, 0;0, 0, 0, 0, 0;0, 0, 0, 0, 0;0, 0, 0, 0, 0]
