OpenCV CUDA模块设备层-----线程块级别的一个内存填充工具函数blockFill()

• 操作系统：ubuntu22.04
• OpenCV版本：OpenCV4.9
• IDE:Visual Studio Code
• 编程语言：C++11

算法描述

在同一个线程块（thread block）内，将 [beg, end) 范围内的数据并行地填充为指定值 value。
它使用了 CUDA 线程协作机制（warp-level 或 block-level） 来实现高效的块级填充，通常比简单的逐线程填充更快。

函数原型

__device__ static __forceinline__ void cv::cudev::blockFill 	
(
 	It  	beg,
	It  	end,
	const T &  	value 
) 		

参数

使用场景

• 初始化图像局部区域；
• 构建 GPU 并行算法前的数据预处理；
• 图像掩码（mask）、ROI 区域清零；
• 需要多个线程协同完成连续内存初始化的任务。

代码示例

#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

// CUDA 核函数：使用 blockFill 填充图像局部区域为指定值
__global__ void fillTileKernel(uchar* data, size_t pitch, int width, int height, uchar value) {
    // 获取当前线程块负责的图像区域起点
    int x = blockIdx.x * TILE_SIZE + threadIdx.x;
    int y = blockIdx.y * TILE_SIZE + threadIdx.y;

    // 检查是否在图像范围内
    if (x >= width || y >= height)
        return;

    // 计算该线程对应的全局指针位置
    uchar* ptr = data + y * pitch + x;

    // 使用 blockFill 填充当前 tile 区域为指定值
    blockFill(ptr, ptr + TILE_SIZE * TILE_SIZE, value);
}

int main() {
    // 创建一个测试图像（8UC1）
    cv::Mat h_src = cv::Mat::zeros(128, 128, CV_8UC1);  // 全黑图像
    std::cout << "Original image sample:\n" << h_src(cv::Rect(0, 0, 5, 5)) << std::endl;

    // 上传到 GPU
    cv::cuda::GpuMat d_src;
    d_src.upload(h_src);

    // 设置 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);

    // 调用核函数，将每个 tile 填充为 255（白色）
    fillTileKernel<<<blocks, threads>>>(d_src.ptr<uchar>(), d_src.step, d_src.cols, d_src.rows, 255);

    // 下载结果
    cv::Mat h_dst;
    d_src.download(h_dst);

    // 显示部分结果
    std::cout << "Filled image sample:\n" << h_dst(cv::Rect(0, 0, 5, 5)) << std::endl;

    return 0;
}

运行结果

Original image sample:
[  0,   0,   0,   0,   0;
   0,   0,   0,   0,   0;
   0,   0,   0,   0,   0;
   0,   0,   0,   0,   0;
   0,   0,   0,   0,   0]
Filled image sample:
[255,   0, 255,   0, 255;
   0,   0,   0,   0,   0;
   0,   0,   0,   0,   0;
   0,   0,   0,   0,   0;
   0,   0,   0,   0,   0]