• 操作系統：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 或類似

主要功能

功能	描述
封裝紋理對象	包含 cudaTextureObject_t
支持 ROI 偏移	提供 (xoff, yoff) 偏移信息
核函數參數傳遞	可作為泛型參數傳入模板核函數
提供采樣接口	使用 tex(y, x) 接口讀取像素值

常用構造函數

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

參數	類型	描述
tex_	cudaTextureObject_t	已創建好的 CUDA 紋理對象
yoff_	int	Y 方向偏移（ROI 起始行）
xoff_	int	X 方向偏移（ROI 起始列）

使用流程總結

步驟	內容
1. 創建 CUDA Array	cudaMallocArray() + cudaMemcpy2DToArray()
2. 配置資源描述符	cudaResourceDesc
3. 配置紋理描述符	cudaTextureDesc
4. 創建紋理對象	cudaCreateTextureObject()
5. 構造 TextureOffPtr	TextureOffPtr<uchar>(texObj, dy, dx)
6. 傳入核函數	使用模板泛型 template <typename TexPtr>
7. 采樣數據	在核函數中使用 tex(y, x)

優點與適用場景

優點	說明
輕量級封裝	不影響性能
支持偏移	可處理 ROI 圖像
易于集成	可作為泛型參數傳入核函數
跨版本兼容	支持 OpenCV ≥ 4.6 所有版本

代碼示例

#include <opencv2/opencv.hpp>
#include <opencv2/cudaimgproc.hpp>
#include <opencv2/cudev/ptr2d/texture.hpp> // TextureOffPtrusing 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;}// 上傳到 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());// 調用基于 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;
}

運行結果