OpenCV CUDA模块设备层-----设备端（GPU）线程块级别的一个内存拷贝工具函数blockCopy()

• 操作系统：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 
) 		

参数

使用场景

• 图像处理中的 局部块拷贝
• 构建 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

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

    // 调用 kernel
    processImageKernel<<<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]