- 操作系統：ubuntu22.04
- OpenCV版本：OpenCV4.9
- IDE:Visual Studio Code
- 編程語言：C++11
算法描述
cv::cudev::Texture 是 OpenCV CUDA 模塊（opencv_cudaimgproc）中用于 CUDA 紋理內存（Texture Memory） 的封裝類。它主要用于在 CUDA 核函數中訪問圖像數據時，利用紋理內存的緩存機制來提升性能，特別是在圖像采樣、縮放、仿射變換等操作中。
- 紋理內存是一種 只讀緩存內存，適合隨機訪問模式。
- 它有緩存機制，對圖像空間局部性訪問非常友好。
- 常用于圖像處理中的插值、旋轉、透視變換、濾波等任務。
cv::cudev::Texture 的作用
OpenCV 提供了 cv::cudev::Texture 類模板來綁定圖像數據到紋理內存中：
template <typename T>
class Texture : public PtrStepSzb
{
public:Texture();explicit Texture(const GpuMat& d_src);void bind(const GpuMat& d_src);void unbind();
};
你可以把它理解為一個“GPU 圖像紋理對象”，綁定后可以在核函數中使用類似 CPU 中 cv::getRectSubPix 或 cv::remap 的方式訪問像素。
使用步驟示例
步驟 1：包含頭文件
#include <opencv2/cudaimgproc.hpp>
#include <opencv2/cudev/ptr2d/texture.hpp>
步驟 2：綁定圖像到紋理內存
cv::cuda::GpuMat d_src = ...; // 輸入圖像
cv::cudev::Texture<uchar> tex;
tex.bind(d_src); // 綁定 uchar 類型的圖像到紋理內存
你也可以使用其他類型如 uchar3, float 等。
步驟 3：在 CUDA 核函數中訪問紋理內存
__global__ void sampleKernel(float* output, int width, int height) {int x = blockIdx.x * blockDim.x + threadIdx.x;int y = blockIdx.y * blockDim.y + threadIdx.y;if (x < width && y < height) {float u = static_cast<float>(x) / width;float v = static_cast<float>(y) / height;// 使用紋理采樣器讀取像素output[y * width + x] = tex2D(tex, u * width, v * height);}
}
注意：這里使用了 tex2D() 函數，這是 CUDA 運行時 API 中的標準紋理采樣函數。
步驟 4：調用核函數并釋放資源
dim3 block(16, 16);
dim3 grid((width + 15) / 16, (height + 15) / 16);sampleKernel<<<grid, block>>>();
cudaDeviceSynchronize();tex.unbind(); // 使用完記得解綁
注意事項
| 內容 | 說明 |
|---|---|
| 只讀訪問 | Texture memory 是只讀的，不能寫入 |
| 數據類型支持 | 支持 uchar, uchar4, float, float4 等常見格式 |
| 性能優化 | 對圖像縮放、旋轉、仿射變換等操作特別有用 |
| 自動邊界處理 | 支持 cudaAddressModeClamp, cudaAddressModeWrap 等尋址模式 |
| 需要綁定和解綁 | 使用完畢要調用 unbind() 避免資源泄漏 |
代碼示例
頭文件：
#ifndef CUDA_UTILS_H
#define CUDA_UTILS_H#include <opencv2/opencv.hpp>
#include <opencv2/cudaimgproc.hpp>// 聲明在 .cu 文件中實現的函數
void resizeWithTexture(cv::cuda::GpuMat& d_src, cv::cuda::GpuMat& d_dst, float scale);#endif // CUDA_UTILS_H‘
cu文件：
#include "cuda_utils.h"
#include <opencv2/cudev/ptr2d/texture.hpp>using namespace cv;
using namespace cudev;#include <cuda_runtime.h>
#include <vector_types.h>
#include <iostream>// 定義 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, int src_cols, int src_rows, cudaTextureObject_t texObj) {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;// 使用紋理采樣器讀取像素值float val = tex2D<uchar>(texObj, src_x + 0.5f, src_y + 0.5f);dst[y * dst_step + x] = static_cast<uchar>(val);}
}void resizeWithTexture(cuda::GpuMat& d_src, cuda::GpuMat& d_dst, float scale) {cudaTextureObject_t texObj = 0;// 1. 創建 CUDA ArraycudaArray* cu_array = nullptr;// 獲取源圖像的通道數并創建對應格式的通道描述符int num_channels = d_src.channels();cudaChannelFormatDesc channel_desc;switch (num_channels) {case 1: channel_desc = cudaCreateChannelDesc<uchar>(); break;case 3: channel_desc = cudaCreateChannelDesc<uchar3>(); break;case 4: channel_desc = cudaCreateChannelDesc<uchar4>(); break;default:std::cerr << "Unsupported number of channels: " << num_channels << std::endl;exit(EXIT_FAILURE);
}CUDA_CHECK(cudaMallocArray(&cu_array, &channel_desc, d_src.cols, d_src.rows, cudaArrayDefault));// 2. 將圖像數據拷貝到 CUDA ArrayCUDA_CHECK(cudaMemcpy2DToArray(cu_array, 0, 0,d_src.data, d_src.step,d_src.cols, d_src.rows,cudaMemcpyDeviceToDevice));// 3. 配置紋理資源描述符cudaResourceDesc res_desc = {};memset(&res_desc, 0, sizeof(res_desc));res_desc.resType = cudaResourceTypeArray;res_desc.res.array.array = cu_array;// 4. 配置紋理描述符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; // 坐標單位為像素而非 [0,1]// 5. 創建紋理對象CUDA_CHECK(cudaCreateTextureObject(&texObj, &res_desc, &tex_desc, NULL));// 6. 啟動核函數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, d_src.cols, d_src.rows, texObj);CUDA_CHECK(cudaDeviceSynchronize());// 7. 清理資源CUDA_CHECK(cudaDestroyTextureObject(texObj));CUDA_CHECK(cudaFreeArray(cu_array));
}
main.cpp:
#include "cuda_utils.h" // 調用 CUDA 接口
#include <iostream>
#include <opencv2/cudaimgproc.hpp>
#include <opencv2/opencv.hpp>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() );// 調用 CUDA 實現的縮放函數resizeWithTexture( d_src, d_dst, scale );// 下載結果cv::Mat h_dst;d_dst.download( h_dst );// 顯示結果cv::imshow( "Original", h_src );cv::imshow( "Resized (CUDA Texture)", h_dst );cv::waitKey( 0 );return 0;
}
運行結果
)