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://github.com/colmap/colmap
为什么要用到纹理内存
在实现patch match 算法的时候，要获取图片上亚像素级别的亮度值，说白了，就是比如：要获取 (100.2,100.5)处的像素值，一般的操作是使用双线性插值（不清楚的请自行百度）的方式去计算，但是对于CUDA程序来说，其局部内存、寄存器内存和共享内存都很小，把整幅图像拷贝到这些内存中不太现实，如果直接访问全局内存的话，会降低速度，笔者对比了下自己写插值的方法，在1050ti上能够加速30%左右（具体可能因为环境与硬件不同而有差异）。
简单的代码示例
注： 这里只介绍如何同时使用多个纹理内存的，单个纹理内存的应用例子可以直接参考 CUDA C PROGRAMMING GUIDE中的例子，本文只适用于这些纹理图片大小都相等的状况（当然也可以把不同大小的填充成一样大的）。如在patch match中，所有的参考图像都可以当作纹理内存来看待。 代码主要参考colmap中的代码，但是其他链接中的也都进行了参考
目前为止，如果用到很多个纹理的话。本文应该是比较好的方案。
其它的方法反正笔者试了很多都不能用（如传纹理的指针、纹理的数组之类的）
- 定义变量
//cudaTextureType2DLayered可以理解成分层次访问，有很多个二维的纹理组成
texture<float, cudaTextureType2DLayered, cudaReadModeElementType> texref;
//这个东西是与纹理内存一起使用的
cudaArray* cuArray3d;
//这个只是个别名
#define real float
- 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);
//这个就是数据类型
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(¶ms);
自己创建连续的数组的例子
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));
}
- 设置并绑定纹理内存，纹理内存只有和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)
- 核函数和设备函数中的访问，将之前定义的texture变量定义成全局的就可以了，这个是可以直接在核函数和设备函数中被调用的
切记加上0.5，虽然有的时候不加也是正确的
real value = tex2DLayered(texref, gidx+0.5, gidy+0.5, gidz);
总结
- 纹理内存和cuArray是要一起使用的
- 使用cudaTextureType2DLayered可以方便地访问多个纹理，但是初始化方式比较复杂，数组的绑定要按照文中的方式
- 使用tex2DLayered获取值的时候一定要加上0.5
