CUDA技术帖： 3D texture linear data VS 3D cudaArray

众所周知，Global memory没有Cache，访问速度很慢，Shared memory访问速度很快，但是容量很小，对于较大的数组，将其绑定至texture memory往往是个不错的选择。Texture memory可以cache，而且容量很大。  

在当前的CUDA版本中，3D的线性内存是无法直接绑定到texture memory，一维的可以，因此，需要将数据首先放进一个3D的CUDA array，然后将3D CUDA array绑定到texture memory上，访问数组元素时，通过取纹理的函数tex3D(tex,x,y,z)可以返回坐标为（x,y,z）的元素。  

1. 创建CUDA 3D array  
    在之前的CUDA版本中，extent.width与height，depth不同，其计数单位为bytes，所以在旧版本中必须使用array_width*sizeof(float)，最新的3.1竟然悄悄的修改了。可以CUDA的文档一直是错误的，文档中记载width,height,depth均是in bytes，实际上赋值时使用元素个数即可。如果不直接赋值，还可以调用函数make_cudaExtent(extent,width,height,depth), 原理类似。  

       cudaArray *d_u  
       cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();    
       cudaExtent extent;       
       extent.width=array_width;       
       extent.height=array_height;       
       extent.depth=array_depth;        
       cudaMalloc3DArray(&d_u,&channelDesc,extent);  

2. 复制数据至3D array  
     首先解释一下pitched pointer的工具原理，如果访问数组元素u[x][y][z]，通过pitched pointer访问则是u_p[x+y*pitch+ z*pitch*height ]。 显然，这里pitch=width，因此当创建pitched pointer时我们需要将width和height作为参数传递给函数make_cudaPitchedPtr()。在这里尤其要注意的是，pitched pointer指向的array与传统的C语言数组的存储方式不同，C语言访问元素u[x][y][z]是通过u[x*height*depth+y*depth+z]。因此为了正确读取所需元素，我建议逆序建立pitched pointer：  
       copyParams.srcPtr   = make_cudaPitchedPtr((void*)u, array_depth*sizeof(float), array_depth, array_height);  
此时相当于数组u[x][y][z]被转置，在CUDA3D array中对应元素为u[z][y][x]，CUDA文档与指南中并未提及这一点区别，这个问题当时也困扰我很久，费尽周折才搞清楚，希望以后的SDK sample能覆盖这个注意点。  

       cudaMemcpy3DParms copyParams = {0};     
       copyParams.srcPtr   = make_cudaPitchedPtr((void*)u, array_width*sizeof(float), array_width, array_height);  
       copyParams.dstArray = d_u;  
       copyParams.extent   = extent;  
       copyParams.kind     = cudaMemcpyHostToDevice;  
       cudaMemcpy3D(&copyParams);  

3. 绑定3D array至texture memory  

       normalized 设置是否对纹理坐标是否进行归一化。如果normalized是一个非零值，那么就会使用归一化到[0，1)的坐标进行寻址，否则对尺寸为width, height, depth的纹理使用坐标[0,width-1], [0,height-1], [0,depth-1]寻址。例如，一个尺寸为64×32的纹理可以通过x维度范围为[0，63]，y维度范围[0,31]的坐标寻址。如果采用归一化方式对尺寸为64×32的纹理进行寻址，在x和y维度上的坐标就都是[0.0,1.0)。这样就可以保证纹理的坐标与纹理的尺寸无关。  
        filterMode用于设置纹理的滤波模式，即如何根据坐标计算返回的纹理值。滤波模式可以是cudaFilterModePoint或者cudaFilterModeLinear。滤波模式为CudaFilterModePoint时，返回值是与坐标最接近的像元的值。CudaFilterModeLinear模式只能对返回值为浮点型的纹理使用，启用这一种模式时将拾取纹理坐标周围的像元，然后根据坐标与这些像元之间的距离进行插值计算。对一维纹理可以使用线性滤波，对二维纹理可以使用双线性滤波。返回值会是对最接近纹理坐标的两个像元（对一维纹理），四个像元（对二维纹理）或者八个像元（对三维纹理）进行插值后得到的值。  

       texture<float,3,cudaReadModeElementType> tex_u;  
       tex_u.filterMode = cudaFilterModePoint;  
       tex_u.normalized = false;  
       tex_u.channelDesc = channelDesc;  
       if (cudaBindTextureToArray(tex_u, d_u, channelDesc) != (unsigned int) CUDA_SUCCESS) {  
         printf("[ERROR] Could not bind texture u\n");  
         return;  
       }