- 操作系统：ubuntu22.04
- OpenCV版本：OpenCV4.9
- IDE:Visual Studio Code
- 编程语言：C++11
算法描述
cv::cudev::TextureOff<T, P> 是 OpenCV 的 CUDA 模块（opencv_cudev） 中用于在 GPU 上使用纹理内存的一个类模板，它主要用于将图像数据绑定到 CUDA 纹理对象，并支持偏移量访问。
定义：
template <typename T, typename Ptr>
class TextureOff : public Texture<T>
- T：纹理中元素的类型（如 uchar, float）
- Ptr：指向图像数据的指针类型（如 GlobPtrSz）
这个类封装了从图像数据创建 CUDA 约束纹理对象的过程，并支持带有 ROI（Region of Interest）偏移量的访问。
主要功能
- 自动将图像数据上传到 CUDA Array
- 创建并管理 cudaTextureObject_t
- 支持带偏移量的纹理访问（offsetX(), offsetY()）
- 可以像普通纹理一样在核函数中使用
相关类和结构体
| 类/结构体 | 作用 |
|---|---|
| cv::cudev::Texture | 纹理基类 |
| cv::cudev::TextureOff<T, P> | 支持偏移的纹理类 |
| cv::cudev::GlobPtrSz | 封装 GpuMat 数据的指针包装器 |
| cudaTextureObject_t | CUDA 原生纹理对象句柄 |
| cudaResourceDesc | 描述纹理资源来源（如 array） |
| cudaTextureDesc | 描述纹理采样方式（filter mode, address mode 等） |
代码示例
注意：是cu文件
#include <opencv2/opencv.hpp>
#include <opencv2/cudaimgproc.hpp>#include <iostream>using namespace cv;
using namespace cuda;// 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)// 核函数：使用纹理对象进行图像缩放
__global__ void resizeKernel(uchar* dst, int dst_cols, int dst_rows, size_t dst_step,float scale, cudaTextureObject_t texObj, int dx, int dy) {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) + dx;float src_y = (y / scale) + dy;// 使用纹理对象采样uchar val = tex2D<uchar>(texObj, src_x + 0.5f, src_y + 0.5f);dst[y * dst_step + x] = val;}
}void resizeWithTextureOff(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));// 启动核函数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>>>(d_dst.data, d_dst.cols, d_dst.rows, d_dst.step,scale, texObj, 0, 0);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;}// 上传到 GPUcv::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());// 调用基于 Texture 的缩放函数resizeWithTextureOff(d_src, d_dst, scale);// 下载结果cv::Mat h_dst;d_dst.download(h_dst);// 归一化显示cv::Mat h_dst_normalized;cv::normalize(h_dst, h_dst_normalized, 0, 255, cv::NORM_MINMAX, CV_8UC1);// 显示图像cv::imshow("Original", h_src);cv::imshow("Resized (Texture)", h_dst_normalized);cv::waitKey(0);return 0;
}
运行结果
