CUDA纹理内存--硬件插值功能的应用
1. 纹理内存的使用方式
纹理内存是CUDA的一种只读内存，通常的使用方式有两种：
(1) 把数据从host端拷贝到device端的CUDA数据，然后将CUDA数组绑定到纹理内存，通过访问纹理内存（也称为纹理拾取）来获取CUDA数组中的数据。
(2) 把数据从host端拷贝到device端某一段连续的全局内存，然后将该段全局内存绑定到纹理内存，通过纹理拾取来获取全局内存中的数据。
2. 使用纹理内存的优势
(1) 纹理内存针对二维数据的局部访问作了优化，因此通过纹理拾取访问二维数据的速度，相比于直接访问设备全局内存更快。
所以如果在核函数中需要频繁访问全局内存，可以考虑纹理拾取的方式。
(2) 纹理内存具有硬件插值功能，包括最邻近插值和双线性插值这两种插值方式。如果纹理拾取时输入的访问坐标地址是浮点数，纹理内存将自动根据设置插值方式对浮点坐标进行插值，然后返回插值结果。这个插值过程不需要开发者来实现，是硬件自动完成的，开发者只需要设置好插值方式为最邻近插值或者双线性插值即可，因此可以节省很多计算时间。
因此在CUDA优化程序中如果需要浮点坐标插值，可以考虑使用纹理内存的硬件插值特性来加速计算。不过需要注意，其仅当纹理内存的数据为float浮点型数据时，双线性插值模式才是有效的。这是本人的亲身经历，一开始把纹理内存设置为uchar型，纹理拾取时输出为空，摸不着头脑，纠结了蛮久才发现是这个原因。
3. 硬件插值功能的应用例子
下面我们使用CUDA来优化图像缩放函数，并在其中使用纹理内存的硬件插值。希望能够通过这个简单的应用例子把硬件插值功能的应用场景与应用方法阐述清楚。
下面上代码，我们在代码后面加上详细注释。
核函数：
//定义纹理内存变量
texture<float, cudaTextureType2D, cudaReadModeElementType> tex_src;
__global__ void resize_img_ker(int row, int col, float x_a, float y_a, uchar *out)
{
int x = threadIdx.x + blockDim.x * blockIdx.x; //col
int y = threadIdx.y + blockDim.y * blockIdx.y; //row
if (x < col && y < row)
{
float xx = x*x_a;
float yy = y*y_a;
//这里的xx和yy都是浮点数，tex2D函数返回的数值已经过硬件插值了，所以不需要开发者再进行插值啦~
out[y*col+x] = (uchar)tex2D(tex_src, xx, yy);
}
}
主体函数：
void resize_img_cuda(Mat src, Mat &dst, float row_m, float col_m)
{
const int row = (int)(src.rows*row_m);
const int col = (int)(src.cols*col_m);
const int srcimg_size = src.rows*src.cols*sizeof(float);
const int dstimg_size = row*col;
const float x_a = 1.0 / col_m;
const float y_a = 1.0 / row_m;
uchar *dst_cuda;
cudaMalloc((void**)&dst_cuda, dstimg_size);
Mat src_tmp;
src.convertTo(src_tmp, CV_32F); //注意这里要把图像转换为float浮点型，否则线性插值模式无法使用
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();//声明数据类型
cudaArray *cuArray_src; //定义CUDA数组
cudaMallocArray(&cuArray_src, &channelDesc, src_tmp.cols, src_tmp.rows); //分配大小为col*row的CUDA数组
tex_src.addressMode[0] = cudaAddressModeWrap;//寻址方式
tex_src.addressMode[1] = cudaAddressModeWrap;//寻址方式 如果是三维数组则设置texRef.addressMode[2]
tex_src.normalized = false;//是否对纹理坐标归一化
tex_src.filterMode = cudaFilterModeLinear;//硬件插值方式：最邻近插值--cudaFilterModePoint 双线性插值--cudaFilterModeLinear
cudaBindTextureToArray(&tex_src, cuArray_src, &channelDesc); //把CUDA数组绑定到纹理内存
cudaMemcpyToArray(cuArray_src, 0, 0, src_tmp.data, srcimg_size, cudaMemcpyHostToDevice); //把源图像数据拷贝到CUDA数组
dim3 Block_resize(16, 16);
dim3 Grid_resize((col + 15) / 16, (row + 15) / 16);
//调用核函数
resize_img_ker << <Grid_resize, Block_resize >> > (row, col, x_a, y_a, dst_cuda);
dst = Mat::zeros(row, col, CV_8UC1);
cudaMemcpy(dst.data, dst_cuda, dstimg_size, cudaMemcpyDeviceToHost);
cudaFree(dst_cuda);
cudaFreeArray(cuArray_src);
cudaUnbindTexture(tex_src);
}
我们注意到以上代码中，关键是把图像数据转换为float型数据，双线性插值模式（也即cudaFilterModeLinear模式）才可以使用，否则只能使用最邻近插值模式（也即cudaFilterModePoint模式）。我们使用以上实现的函数对248*236的Lena图像的宽和高都放大到原来的五倍，分别设置为cudaFilterModePoint和cudaFilterModeLinear模式，记录运行耗时分别为0.9292ms和0.987ms，而CUDA优化之前，两种插值方式的耗时分别为1.61ms和10.71ms。可以看到加速效果还是比较明显的，特别是双线性插值。
原图
最邻近插值
双线性插值
微信公众号如下，欢迎扫码关注，欢迎私信技术交流：
