CUDA cudaTextureType2DLayered 纹理内存使用心得（踩坑记录）

写本文的原因

笔者是做三维重建的，主要是NVS部分。在实现patch match方式的时候，因需要用到GPU加速，所以需要用到纹理内存来辅助计算。在实现的过程中，踩了不少坑，而且网上的教程非常的少，所以特此记录下来，方便有需要的朋友们参考。因笔者能力有限，所以难免出现很多错误，如发现，希望能够指正。

所参考的资料的链接

1.https://forums.developer.nvidia.com/t/guide-cudamalloc3d-and-cudaarrays/23421
2. https://blog.csdn.net/fb_help/article/details/79711389
3. https://blog.csdn.net/u011291667/article/details/104226734/
4.https://stackoverflow.com/questions/10643790/texture-memory-tex2d-basics
5. https://github.com/colmap/colmap

为什么要用到纹理内存

在实现patch match 算法的时候，要获取图片上亚像素级别的亮度值，说白了，就是比如：要获取 (100.2,100.5)处的像素值，一般的操作是使用双线性插值（不清楚的请自行百度）的方式去计算，但是对于CUDA程序来说，其局部内存、寄存器内存和共享内存都很小，把整幅图像拷贝到这些内存中不太现实，如果直接访问全局内存的话，会降低速度，笔者对比了下自己写插值的方法，在1050ti上能够加速30%左右（具体可能因为环境与硬件不同而有差异）。

简单的代码示例

注： 这里只介绍如何同时使用多个纹理内存的，单个纹理内存的应用例子可以直接参考 CUDA C PROGRAMMING GUIDE中的例子，本文只适用于这些纹理图片大小都相等的状况（当然也可以把不同大小的填充成一样大的）。如在patch match中，所有的参考图像都可以当作纹理内存来看待。 代码主要参考colmap中的代码，但是其他链接中的也都进行了参考
目前为止，如果用到很多个纹理的话。本文应该是比较好的方案。
其它的方法反正笔者试了很多都不能用（如传纹理的指针、纹理的数组之类的）

1. 定义变量

//cudaTextureType2DLayered可以理解成分层次访问，有很多个二维的纹理组成
texture<float, cudaTextureType2DLayered, cudaReadModeElementType> texref;
//这个东西是与纹理内存一起使用的
cudaArray* cuArray3d;
//这个只是个别名
#define real float

2. cudaArray的初始化与拷贝，对于make_cudaPitchedPtr的使用是要特别说明一下的，该函数的参数如下，看上去并不太容易理解。简单的说就是这样之后的访问方式就是 对于“三维数组” img[x][y][z] 的读取为 x + 数组列数 × 数据大小 × y + 数组列数 × 数据大小 × 数组行数 × z。
   把三维数组引起来是因为一般存储的时候都是以一维数组存储的，就比如对于两张图片，先将第一张图片按照行依次写入，然后再将第二幅图片按行依次写入。这里写的比较简洁，更详细的可参考链接2

static __inline__ __host__ struct cudaPitchedPtr make_cudaPitchedPtr(void *d, size_t p, size_t 
xsz, size_t ysz);
/*
param d   - Pointer to allocated memory
param p   - Pitch of allocated memory in bytes
param xsz - Logical width of allocation in elements
param ysz - Logical height of allocation in elements
*/

//这个就是数据类型
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<real>();
//cudaExtent 是一个结构体，在这里depth相当于纹理图片的数量，长宽自己根据实际情况定义
cudaExtent extent;
extent.width = maxWidth;
extent.height = maxHeight;
extent.depth = neighborNum;

//申请下内存	
cudaMalloc3DArray(&cuArray3d, &channelDesc, extent, cudaArrayLayered);

//内存拷贝
cudaMemcpy3DParms params = { 0 };
params.extent = extent;
params.kind = cudaMemcpyHostToDevice;
params.dstArray = cuArray3d;
//就如之前所说的那样，应该不太难理解，注意neighborImgData一般需要自己写一个，最好按照先行后列再图片的方式写，下面会给一个例子
params.srcPtr =make_cudaPitchedPtr((void*)neighborImgData, maxWidth * sizeof(real), maxWidth, maxHeight);
cudaMemcpy3D(&params);

自己创建连续的数组的例子

int maxImgMemSize = maxHeight * maxWidth * sizeof(real);
real* imgData =(real*) malloc(maxImgMemSize*imgNum);
real* imgDataTemp =imgData;
//imgSet[i]为 cv::Mat 组成的vector
for(int i = 0; i < imgSize; i++){
	//如果使用opencv的时候需要填充边缘可以用以下的函数填充，参数自己看
	//cv::copyMakeBorder(img,img,0,10,0,10,cv::BORDER_CONSTANT,cv::Scalar(0.0));
	memcpy(imgDataTemp , imgSet[i].data, maxImgMemSize);
	imgDataTemp += (maxImgMemSize / sizeof(real));
}

3. 设置并绑定纹理内存，纹理内存只有和cuArray绑定才能用

//这两个也就是在对于超出了边界的寻址的时候如何处理，本次采用的是按照0来处理
texref.addressMode[0] = cudaAddressModeBorder;
texref.addressMode[1] = cudaAddressModeBorder;
//这个属于访问方式，可以将访问的范围限定到0-1，这里用不到，就设定成false
texref.normalized = false;
//使用线性插值来计算，使用这个选项的时候一定要注意，一般是向下取整，所以要加0.5
//避免无法精确表示的某些数值出现错误取值，如 x=3，实际是2.99999，此时实际获取的是x=2的元素
texref.filterMode = cudaFilterModeLinear;
cudaBindTextureToArray(neighborImgTexture, cuArray3d);
cudaFreeArray(cuArray3d)

4. 核函数和设备函数中的访问，将之前定义的texture变量定义成全局的就可以了，这个是可以直接在核函数和设备函数中被调用的
   切记加上0.5，虽然有的时候不加也是正确的
   对于这个问题，链接3和4的可以一起参考，对于不同的插值方式+0.5的意义是不同的，英文好的话建议直接看stackoverflow的那个

real value = tex2DLayered(texref, gidx+0.5, gidy+0.5, gidz);

总结

1. 纹理内存和cuArray是要一起使用的
2. 使用cudaTextureType2DLayered可以方便地访问多个纹理，但是初始化方式比较复杂，数组的绑定要按照文中的方式
3. 使用tex2DLayered获取值的时候一定要加上0.5