如何在CUDA核函数中正确使用OpenCV Mat并基于GpuMat重构代码

1. OpenCV Mat无法在设备端使用：Mat是主机端专属的数据结构，其成员函数（如at<>）仅编译为主机代码，CUDA核函数（设备端）无法调用这些函数。
2. 错误的内存拷贝逻辑：直接拷贝Mat对象到设备毫无意义——Mat本身只是包含元数据的结构体，真正的图像数据在主机内存的data指针指向区域，设备端的Mat结构体无法访问主机内存。

以下提供两种重构方案，优先推荐使用OpenCV官方的GpuMat简化设备内存管理。

方案一：使用OpenCV GpuMat重构 #

GpuMat是OpenCV专为CUDA设计的设备端矩阵结构，自动管理设备内存，无需手动调用cudaMalloc/cudaFree。

主机函数重构 #

#include <opencv2/core/cuda.hpp>

void myBlur(Mat face, int w, int h, int THREADS, int r) {
    Mat table = summed_table(face, w, h);
    int blur_w = w - r - 1;
    int blur_h = h - r - 1;

    // 1. 主机Mat转设备GpuMat
    cv::cuda::GpuMat d_face(face);
    cv::cuda::GpuMat d_table(table);
    cv::cuda::GpuMat d_result(blur_h, blur_w, face.type()); // 预分配结果内存

    // 2. 配置CUDA线程网格（向上取整避免遗漏边缘像素）
    dim3 threadsPerBlock(THREADS, THREADS);
    dim3 numBlocks((blur_w + threadsPerBlock.x - 1) / threadsPerBlock.x,
                   (blur_h + threadsPerBlock.y - 1) / threadsPerBlock.y);

    // 3. 调用核函数，传入设备指针、步长等关键参数
    myBlurKernel<<<numBlocks, threadsPerBlock>>>(
        reinterpret_cast<Vec3i*>(d_result.data), d_result.step,
        reinterpret_cast<const Vec3i*>(d_table.data), d_table.step,
        blur_w, blur_h, r
    );

    // 4. 设备结果拷贝回主机Mat（仅覆盖有效模糊区域）
    d_result.download(face(cv::Rect(0, 0, blur_w, blur_h)));
}

核函数重构 #

__global__ void myBlurKernel(
    Vec3i* d_result, size_t result_step,
    const Vec3i* d_table, size_t table_step,
    int width, int height, int r
) {
    // 修正线程索引：blockIdx.y对应行，blockIdx.x对应列
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;

    const int area = (2 * r + 1) * (2 * r + 1); // 用整数乘法替代pow，避免浮点开销

    if (row < height && col < width) {
        // 计算每行的元素步长（将字节步长转为Vec3i数量）
        size_t table_row_stride = table_step / sizeof(Vec3i);
        size_t result_row_stride = result_step / sizeof(Vec3i);

        // 定位积分表的四个角点
        const Vec3i* top_right = d_table + (row + r) * table_row_stride + (col + r);
        const Vec3i* top_left = d_table + (row + r) * table_row_stride + (col - r - 1);
        const Vec3i* bottom_right = d_table + (row - r - 1) * table_row_stride + (col + r);
        const Vec3i* bottom_left = d_table + (row - r - 1) * table_row_stride + (col - r - 1);

        // 计算区域求和并平均
        Vec3i sum;
        sum[0] = (top_right[0] - top_left[0] - bottom_right[0] + bottom_left[0]) / area;
        sum[1] = (top_right[1] - top_left[1] - bottom_right[1] + bottom_left[1]) / area;
        sum[2] = (top_right[2] - top_left[2] - bottom_right[2] + bottom_left[2]) / area;

        // 写入结果
        d_result[row * result_row_stride + col] = sum;
    }
}

方案二：手动管理设备内存（不使用GpuMat） #

若不想依赖GpuMat，可直接操作设备内存指针，但需手动处理数据拷贝和内存释放。

主机函数重构 #

void myBlur(Mat face, int w, int h, int THREADS, int r) {
    Mat table = summed_table(face, w, h);
    int blur_w = w - r - 1;
    int blur_h = h - r - 1;

    size_t face_data_size = face.total() * face.elemSize();
    size_t table_data_size = table.total() * table.elemSize();
    size_t result_data_size = blur_h * blur_w * sizeof(Vec3i);

    // 分配设备内存
    Vec3i* d_face_data = nullptr;
    Vec3i* d_table_data = nullptr;
    Vec3i* d_result_data = nullptr;
    cudaMalloc(&d_face_data, face_data_size);
    cudaMalloc(&d_table_data, table_data_size);
    cudaMalloc(&d_result_data, result_data_size);

    // 主机数据拷贝到设备
    cudaMemcpy(d_face_data, face.data, face_data_size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_table_data, table.data, table_data_size, cudaMemcpyHostToDevice);

    // 配置线程网格
    dim3 threadsPerBlock(THREADS, THREADS);
    dim3 numBlocks((blur_w + threadsPerBlock.x - 1) / threadsPerBlock.x,
                   (blur_h + threadsPerBlock.y - 1) / threadsPerBlock.y);

    // 调用核函数
    myBlurKernel<<<numBlocks, threadsPerBlock>>>(
        d_result_data, blur_w * sizeof(Vec3i),
        d_table_data, table.step,
        blur_w, blur_h, r
    );

    // 结果拷贝回主机
    cudaMemcpy(face.data, d_result_data, result_data_size, cudaMemcpyDeviceToHost);

    // 释放设备内存
    cudaFree(d_face_data);
    cudaFree(d_table_data);
    cudaFree(d_result_data);
}

核函数与方案一完全一致 #

• 确保summed_table返回的积分表类型为CV_32SC3，与输入图像类型匹配，否则会出现计算错误。
• 建议将r作为参数传入主机函数，避免依赖全局变量。
• 调试时可添加CUDA错误检查：

  cudaDeviceSynchronize();
  cudaError_t err = cudaGetLastError();
  if (err != cudaSuccess) {
      printf("CUDA错误：%s\n", cudaGetErrorString(err));
  }
  

内容的提问来源于stack exchange，提问作者Santiago Vallejo