OpenCV CUDA模块设备层-----用于封装CUDA纹理对象+ROI偏移量的一个轻量级指针类TextureOffPtr()

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

算法描述

TextureOffPtr<T, R> 是 OpenCV 的 CUDA 模块（opencv_cudev）中用于封装 CUDA 纹理对象 + ROI 偏移量 的一个轻量级指针类。它允许在核函数中像访问普通纹理一样进行采样，同时自动考虑图像的偏移位置。

• T：纹理元素类型（如 uchar, float）
• R：底层指针类型，默认为 Texture2DLayeredPtr 或类似

主要功能

常用构造函数

__host__ TextureOffPtr(const cudaTextureObject_t tex_, const int yoff_, const int xoff_)

使用流程总结

优点与适用场景

代码示例

#include <opencv2/opencv.hpp>
#include <opencv2/cudaimgproc.hpp>
#include <opencv2/cudev/ptr2d/texture.hpp> // TextureOffPtr

using namespace cv;
using namespace cudev;

// CUDA 错误检查宏
#define CUDA_CHECK(call) \
    do { \
        cudaError_t err = call; \
        if (err != cudaSuccess) { \
            std::cerr << "CUDA error at " << __FILE__ << ":" << __LINE__ << ": " \
                      << cudaGetErrorString(err) << std::endl; \
            exit(EXIT_FAILURE); \
        } \
    } while (0)

// 核函数：使用 TextureOffPtr 进行图像缩放（最近邻插值）
template <typename TexPtr>
__global__ void resizeKernel(TexPtr tex, uchar* dst, int dst_cols, int dst_rows, size_t dst_step, float scale) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    if (x < dst_cols && y < dst_rows) {
        float src_x = x / scale;
        float src_y = y / scale;
        dst[y * dst_step + x] = tex(src_y, src_x); // 自动加偏移
    }
}

void resizeWithTextureOffPtr(cuda::GpuMat& d_src, cuda::GpuMat& d_dst, float scale) {
    int width = d_src.cols;
    int height = d_src.rows;

    // 创建 CUDA Array 并拷贝数据
    cudaChannelFormatDesc channel_desc = cudaCreateChannelDesc<uchar>();
    cudaArray* cu_array = nullptr;
    CUDA_CHECK(cudaMallocArray(&cu_array, &channel_desc, width, height));

    CUDA_CHECK(cudaMemcpy2DToArray(
        cu_array, 0, 0,
        d_src.data, d_src.step,
        width, height,
        cudaMemcpyDeviceToDevice));

    // 配置纹理资源描述符
    cudaResourceDesc res_desc = {};
    memset(&res_desc, 0, sizeof(res_desc));
    res_desc.resType = cudaResourceTypeArray;
    res_desc.res.array.array = cu_array;

    // 配置纹理描述符
    cudaTextureDesc tex_desc = {};
    memset(&tex_desc, 0, sizeof(tex_desc));
    tex_desc.addressMode[0] = cudaAddressModeClamp;
    tex_desc.addressMode[1] = cudaAddressModeClamp;
    tex_desc.filterMode = cudaFilterModePoint;
    tex_desc.readMode = cudaReadModeElementType;
    tex_desc.normalizedCoords = 0;

    // 创建纹理对象
    cudaTextureObject_t texObj = 0;
    CUDA_CHECK(cudaCreateTextureObject(&texObj, &res_desc, &tex_desc, NULL));

    // 设置偏移量（假设无 ROI）
    int dx = 0;
    int dy = 0;

    // 使用 TextureOffPtr 封装纹理对象和偏移
    TextureOffPtr<uchar> texPtr(texObj, dy, dx);

    // 启动核函数
    dim3 block(16, 16);
    dim3 grid((d_dst.cols + block.x - 1) / block.x,
              (d_dst.rows + block.y - 1) / block.y);

    resizeKernel<<<grid, block>>>(texPtr, d_dst.data, d_dst.cols, d_dst.rows, d_dst.step, scale);

    CUDA_CHECK(cudaDeviceSynchronize());

    // 清理资源
    CUDA_CHECK(cudaDestroyTextureObject(texObj));
    CUDA_CHECK(cudaFreeArray(cu_array));
}

int main() {
    // 加载图像（灰度图）
    cv::Mat h_src = cv::imread("/media/dingxin/data/study/OpenCV/sources/images/Lenna.png", cv::IMREAD_GRAYSCALE);
    if (h_src.empty()) {
        std::cerr << "Failed to load image!" << std::endl;
        return -1;
    }

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

    // 设置目标尺寸（放大两倍）
    float scale = 2.0f;
    d_dst.create(cvRound(h_src.rows * scale), cvRound(h_src.cols * scale), h_src.type());

    // 调用基于 TextureOffPtr 的缩放函数
    resizeWithTextureOffPtr(d_src, d_dst, scale);

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

    // 显示图像
    cv::imshow("Original", h_src);
    cv::imshow("Resized (TextureOffPtr)", h_dst);
    cv::waitKey(0);

    return 0;
}

运行结果