- 操作系统：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;
}
运行结果
)