cuda三维纹理内存的使用

有时候需要使用cuda处理多帧的图像，需要把多帧时间序列的图像传入到显存中，这个时候就可以把保存图像的全局内存绑定到二维纹理内存，核函数通过纹理拾取来访问输入的图像数据。不过二维纹理内存的宽是有限制的：

  cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
  cudaMallocArray((cudaArray**)&arr_mat_x, &channelDesc, img_size, IIR_N_X); 

在以上定义二维cuda数组的代码中，img_size是有限制的，如果超过64K就会出错。所以如果每帧图像的大小超过了64K，则不能使用二维纹理内存，这时候可以使用三维纹理内存。下面将详细说明三维纹理内存的使用。

1. 定义三维cuda数组：

  cudaArray *arr_mat_x;   //定义device端指针
  cudaExtent extent;   //定义图像的尺寸和帧数结构体
  cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<uchar>();  //定义数据类型
  extent.width = col;    //创建extent时，在旧版本中宽度w以字节为单位，即必须乘上sizeof(DTYPE)，新版本已经不用乘以sizeof(DTYPE)，否则反而会出错!文档和函数说明都没有改过来，这里是坑！
  extent.height = row;   //每帧图像的行数
  extent.depth = picnum;   //图像的总帧数
  cudaMalloc3DArray((cudaArray**)&arr_mat_x, &channelDesc, extent);   //创建picnum帧row*col的存储空间

2. 将数据从CPU拷贝到cuda数组中：

    cudaMemcpy3DParms HostToDev = {0};   //定义数据传输的结构体
    HostToDev.dstArray = arr_mat_x;    //指定数据传输的目标地址为cuda数组
    HostToDev.extent = make_cudaExtent(col, row, 1);      //创建extent时，在旧版本中宽度w以字节为单位，即必须乘上sizeof(DTYPE)，新版本已经不用乘以sizeof(DTYPE)，否则反而会出错!文档和函数说明都没有改过来，这里是坑！  
	HostToDev.kind = cudaMemcpyHostToDevice;  //定义传输方向为CPU到GPU显存
	HostToDev.srcPos = make_cudaPos(0, 0, 0);   //定义数据传输的源地址的偏移量
	for(int i = begin; i < begin+picnum; i++)   //拷贝多帧图像到cuda数组
	{
		HostToDev.srcPtr = make_cudaPitchedPtr((void *)pic[i].data, col*sizeof(uchar), col, row);     //指定数据传输的源地址，注意这里的第二个参数需要乘以数据类型所占的字节数   
		HostToDev.dstPos = make_cudaPos(0, 0, i-begin);   //指定目标地址的偏移量，分别为x，y，z地址
		cudaMemcpy3D(&HostToDev);     //根据以上设置的参数实行拷贝
	}

3. 绑定纹理内存

    texture<uchar, cudaTextureType3D, cudaReadModeElementType>  tex_mat;    //定义为3D类型的纹理内存
	
	tex_mat.normalized = 0;   //不归一化
    tex_mat.filterMode = cudaFilterModePoint;    //filterMode：滤波模式。仅对绑定 CUDA 数组的纹理有效。当使用浮点型的坐标寻址纹理时，将根据设定返回不同类型的值。设定可以有：cudaFilterModePoint 和 cudaFilterModeLinear。分别表示最近邻插值和线性插值
    tex_mat.addressMode[0] = cudaAddressModeClamp;  //寻址模式，即如何处理越界的纹理坐标。可设置：cudaAddressModeClamp 和 cudaAddressModeWrap。Clamp 即钳位模式，Wrap 为循环模式。循环模式只支持归一化的纹理坐标
    tex_mat.addressMode[1] = cudaAddressModeClamp;
    tex_mat.addressMode[2] = cudaAddressModeClamp;
    tex_mat.channelDesc = channelDesc;   //描述纹理返回值类型，同cuda数组部分的内容
	cudaBindTextureToArray(tex_mat, (cudaArray *)arr_mat_x, channelDesc);   //绑定纹理内存

4. 纹理拾取：

tex3D(tex_mat, (float)(j-1), (float)(i-1), (float)k); //后面三个参数分别是x，y，z坐标