声明CUDA数组之前，必须先以结构体channelDesc描述CUDA数组中的数据类型。
struct cudaChannelFormatDesc {
int x, y, z, w;
enum cudaChannelFormatKind f;
};
|
其中，x, y, z和w分别是每个返回值成员的位数，而f是一个枚举变量，可以取一下几个值：
n
cudaChannelFormatKindSigned，如果这些成员是有符号整型；
n
cudaChannelFormatKindUnsigned，如果这些成员是无符号整型；
n
cudaChannelFormatKindFloat，如果这些成员是浮点型；
然后，我们要确定CUDA数组的维度和尺寸。CUDA数组可以通过cudaMalloc3DArray()或cudaMallocArray()函数分配。用cudaMalloc3DArray可以分配一维、二维或者三维的CUDA数组，而cudaMallocArray()一般用于分配二维CUDA数组。在使用完CUDA数组后，要使用cudaFreeArray函数释放显存。
由cudaMalloc3DArray分配的CUDA数组使用cudaMemcpy3D()完成与其他CUDA数组或者线性内存的数据传输。CUDA API中使用结构体cudaExtent描述3D Array和3D线性内存在三个维度上的尺寸，在描述一维、二维和三维数组分别用以下的形式：
cudaextent extent = make_cudaextent([1,8192],0,0);
cudaextent extent = make_cudaextent([1,65535],[1,32768],0);
cudaextent extent = make_cudaextent([1,2048],[1,2048],[1,2048]);
其中方括号[]内为允许的寻址范围。注意到二维CUDA数组的第一个维度的寻址范围大于一维CUDA数组的寻址范围，因此在一维CUDA数组的尺寸不够用时，将二维CUDA数组的第二个维度设为1代替一维CUDA数组，获得更大的寻址范围。
下面是声明一个数据类型为char2型，宽×高×深为64×32×16的CUDA 3D数组，对其初始化，最后释放数组的示例代码：
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(8, 8, 0, 0,cudaChannelFormatKindunsigned); //每个像元由两个char构成
cudaExtent extent = make_cudaextent(64,32,16);//建立cudaExtent结构体，描述CUDA数组的维度和尺寸
cudaArray* cuArray;
cudaMalloc3DArray(&cuArray, &channelDesc, extent); //为cuArray开辟空间
缺
cudaFreeArray(cuArray);
下面则是使用cudaMallocArray声明一个由float型构成，尺寸为64×32的CUDA数组，对其赋值，并最后释放的示例代码：
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0,cudaChannelFormatKindunsigned); //每个像元由一个float构成
cudaArray* cuArray;
cudaMallocArray(&cuArray, &channelDesc, 64, 32); //为cuArray开辟空间
cudaMemcpyToArray(cuArray, 0, 0, h_data, &channelDesc);//第二和第三个参数分别表示在宽度和高度上的偏移量，假设h_data中的数据已经初始化
cudaFreeArray(cuArray);
用于在CUDA数组和主机端或者设备端线性内存，以及在CUDA数组间传输数据的函数还有很多，这些还是还有一些异步调用版本，关于这些函数的具体使用方法请参考CUDA Reference mannual。
3.2.4.1 声明纹理参照系
纹理参照系中的一些属性必须在编译时之前被显示声明。纹理参照系通过一个作用范围为全文件的texture型变量声明：
texture<Type, Dim, ReadMode> texRef;
|
其中，
?
Type 确定了由纹理拾取返回的数据类型；Type可以是B3.1节中描述的任意一种由基本整型或者单精度浮点型组成能的1-，2-或者4-元组向量类型。
?
Dim 确定了纹理参照系的维度，默认为1。
?
ReadMode
可以是cudaReadModeNormalizedFloat或者cudaReadModeElementType。如果ReadMode是cudaReadModeNomalizedFloat，并且Type是16-或者8-bit整型，那么返回的值将是一个浮点数。此时，原来整形的值域会被映射到[0.0，1.0]（对无符号整型），或者[-1.0，1.0]（对有符号整型）。例如，一个值为0xff的8-bit无符号整型会被映射为1.0f。如果使用cudaReadModeElementType,那么就不会对输出进行转换。ReadMode是一个可选参数，如果不写，那么默认就是cudaReadModeElementType。
例如，下面的代码声明了一个二维，像元数据为unsigned char型，但将返回值转换为float型的纹理参照系：
texture<unsigned char, 2, cudaReadModeNormalizedFloat> texRef;
3.2.4.2 设置运行时纹理参照系属性
纹理参照系中的其它属性可以不必声明，并在运行时进行修改。这些参数规定了纹理的寻址模式，是否进行归一化，以及纹理滤波。runtimeAPI拥有底层的C风格和高层的C++风格两种接口。高层API中的texture类型是从底层的textureReference中派生而来的。TextureReference是一个下面的代码描述的结构体。
struct textureReference {
int normalized;
enum cudaTextureFilterMode filterMode;
enum cudaTextureAddressMode addressMode[3];
struct cudaChannelFormatDesc channelDesc;
} |
