CUDA-纹理内存

作者：翟天保Steven
版权声明：著作权归作者所有，商业转载请联系作者获得授权，非商业转载请注明出处

纹理内存是什么？

在 GPU 编程中，纹理内存是一种专门用于二维和三维数据读取的存储方式，特别适合图像处理等应用。通过使用纹理内存，开发者可以利用硬件加速的双线性插值功能，显著提高图像缩放、旋转、平滑等操作的性能。

在这篇文章中，我们将通过一个具体的 CUDA 图像放大案例来深入理解纹理内存的使用。本文的示例展示了如何将图像加载到 GPU 的纹理内存中，利用 CUDA 的cudaTextureObject_t对象进行双线性插值，并通过 CUDA 核函数实现图像放大。我们将详细讨论以下几个要点：

CUDA 纹理内存的基本概念：
- 纹理内存是一种只读内存，专门为频繁读取操作设计，可以在 GPU 上使用硬件加速的插值功能。
- 在图像处理过程中，纹理内存允许开发者利用不同的插值模式（如最近邻和双线性插值）来提高性能和图像质量。
纹理内存的创建与使用：
- 本文示例展示了如何将输入图像转换为浮点型格式并复制到 CUDA 数组，再创建纹理对象以供内核函数使用。
- 通过 cudaCreateTextureObject，将输入图像加载到纹理内存中，简化了图像采样位置的计算，并通过硬件加速的双线性插值提高放大效果。
CUDA 核函数进行图像缩放：
- 我们通过自定义的 scaleKernel 内核函数来完成图像缩放操作，纹理对象提供的双线性插值确保了缩放后的图像质量。
- 内核函数根据输出图像的大小计算对应的采样坐标，利用 tex2D 函数从纹理内存中获取像素值并进行插值处理。
性能提升与纹理内存的优势：
- 纹理内存提供的插值和缓存功能使得图像处理的计算性能得到显著提升，特别是在处理大规模图像时。
- 通过 CUDA 事件计时器，我们能够精确记录内核函数的执行时间，并分析其性能优势。

通过文章案例，你将更好地理解纹理内存的作用和使用方式，同时了解到 CUDA 在图像处理领域的强大性能提升效果。

C++代码

ImageProcessing.cuh

#pragma once
#include <cuda_runtime.h>
#include <iostream>
#include <vector>
#include <opencv2/opencv.hpp>
#include <device_launch_parameters.h>using namespace cv;
using namespace std;#define TILE_WIDTH 16// 图像放大-CPU
cv::Mat enlargeImage_CPU(cv::Mat input, float scale);// 图像放大-GPU
cv::Mat enlargeImage_GPU(cv::Mat input, float scale);

ImageProcessing.cu

#include "ImageProcessing.cuh"// 图像放大-CPU
cv::Mat enlargeImage_CPU(cv::Mat input, float scale)
{cv::Mat result;cv::resize(input, result, cv::Size(0, 0), scale, scale, INTER_LINEAR);return result;
}
// 图像比例变化核函数
__global__ void scaleKernel(cudaTextureObject_t texObj, float* output, int outWidth, int outHeight, float scaleX, float scaleY)
{int x = blockIdx.x * blockDim.x + threadIdx.x;int y = blockIdx.y * blockDim.y + threadIdx.y;if (x < outWidth && y < outHeight) {// 计算输入图像中的采样位置float u = x * scaleX;float v = y * scaleY;// 读取纹理并自动进行双线性插值float pixel = tex2D<float>(texObj, u, v);// 将缩放后的像素值写入输出output[y * outWidth + x] = pixel;}
}
// 图像放大-GPU
cv::Mat enlargeImage_GPU(cv::Mat input, float scale)
{// 定义计时器float spendtime = 0.0f;cudaEvent_t start, end;cudaEventCreate(&start);cudaEventCreate(&end);// 获取图像的宽度和高度int imgWidth = input.cols;int imgHeight = input.rows;// 将图像数据转为 float 类型cv::Mat imgFloat;input.convertTo(imgFloat, CV_32FC1);// 缩放后图像尺寸int outWidth = int(imgWidth * scale);int outHeight = int(imgHeight * scale);// 分配设备内存float* d_outputImage;cudaMalloc(&d_outputImage, outWidth * outHeight * sizeof(float));// 分配并复制输入图像到 CUDA 数组（纹理内存）cudaArray* cuArray;cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();cudaMallocArray(&cuArray, &channelDesc, imgWidth, imgHeight);cudaMemcpyToArray(cuArray, 0, 0, imgFloat.ptr<float>(), imgWidth * imgHeight * sizeof(float), cudaMemcpyHostToDevice);// 定义纹理资源描述符cudaResourceDesc resDesc;memset(&resDesc, 0, sizeof(resDesc));resDesc.resType = cudaResourceTypeArray;resDesc.res.array.array = cuArray;// 定义纹理描述符cudaTextureDesc texDesc;memset(&texDesc, 0, sizeof(texDesc));texDesc.addressMode[0] = cudaAddressModeClamp;  texDesc.addressMode[1] = cudaAddressModeClamp;texDesc.filterMode = cudaFilterModeLinear;   texDesc.readMode = cudaReadModeElementType;     texDesc.normalizedCoords = 0;                // 创建纹理对象cudaTextureObject_t texObj = 0;cudaCreateTextureObject(&texObj, &resDesc, &texDesc, nullptr);// 定义CUDA块和线程数dim3 blockDim(TILE_WIDTH, TILE_WIDTH);dim3 gridDim((outWidth + blockDim.x - 1) / blockDim.x, (outHeight + blockDim.y - 1) / blockDim.y);// 计算输出图像在源图像中的比例关系float scaleX = static_cast<float>(imgWidth) / outWidth; float scaleY = static_cast<float>(imgHeight) / outHeight;// 调用CUDA核函数执行缩放操作cudaEventRecord(start, 0);scaleKernel << <gridDim, blockDim >> > (texObj, d_outputImage, outWidth, outHeight, scaleX, scaleY);cudaDeviceSynchronize();cudaEventRecord(end, 0);cudaEventSynchronize(end);cudaEventElapsedTime(&spendtime, start, end);cout << "scaleKernel spend time:" << spendtime << "ms" << endl;// 将处理后的图像数据从GPU内存复制回主机内存cv::Mat output(outHeight, outWidth, CV_32FC1);cudaMemcpy(output.data, d_outputImage, outWidth * outHeight * sizeof(float), cudaMemcpyDeviceToHost);// 转为灰度图像并保存cv::Mat result8U;output.convertTo(result8U, CV_8UC1);// 清理资源cudaDestroyTextureObject(texObj);cudaFreeArray(cuArray);cudaFree(d_outputImage);return result8U;
}

main.cpp

#include "ImageProcessing.cuh"// 预准备过程
void warmupCUDA()
{float* dummy_data;cudaMalloc((void**)&dummy_data, sizeof(float));cudaFree(dummy_data);
}void main()
{// 预准备warmupCUDA();cout << "enlargeImage test begin." << endl;// 加载cv::Mat src = imread("test pic/test3.jpg", 0);float scale = 5.0f;cout << "scale:" << scale << endl;cout << "size: " << src.cols << " * " << src.rows << endl;// CPU版本clock_t s1, e1;s1 = clock();cv::Mat output1 = enlargeImage_CPU(src, scale);e1 = clock();cout << "CPU time:" << double(e1 - s1)<< "ms" << endl;// GPU版本clock_t s2, e2;s2 = clock();cv::Mat output2 = enlargeImage_GPU(src, scale);e2 = clock();cout << "GPU time:" << double(e2 - s2) << "ms" << endl;// 查看输出cv::Mat test1 = output1.clone();cv::Mat test2 = output2.clone();cout << "enlargeImage test end." << endl;}