谈谈显存的使用分析（二）

二、内存与显存的数据传输类型与使用说明
(1)主机端分页内存的使用方式
对于C语言，主机端的分页内存是用calloc或者malloc函数进行一维空间的分配。而对于二维空间的分配则可以使用下面的函数进行分配：
float **get_matrix_float(int m,int n)
{
int i;
float **a;

a=(float **)calloc(m,sizeof(float *));
for(i=0;i return a;
}
分页内存空间的释放采用free即可，二维的情况则用循环释放。
(2)主机端页锁定内存的使用方式
在CUDA2.2以下，仅提供cudaMallocHost函数用于分配页锁定内存，与C语言函数malloc分配分页内存相对应。而从CUDA2.2开始，页锁定内存增加三种新的类型用于主机多线程的portable，用于高效写回write-combined以及零拷贝的mapping，用cudaHostAlloc进行分配，其中采用四个可选参数标志用以指定使用何种特性：
• cudaHostAllocDefault: 默认情况等价于cudaMallocHost().
• cudaHostAllocPortable: 用于主机多线程数据共享的portable方式
• cudaHostAllocMapped: 通过映射方式分配给CUDA地址空间实现zero-copy，在设备上使用的指针通过调用cudaHostGetDevicePointer()获取
• cudaHostAllocWriteCombined: 分配写联合的存储空间，主要用于主机到设备的传输或者通过映射页锁定空间CPU写而设备读的情况。但CPU读的效率不高。
函数cudaFreeHost()用于释放页锁定空间。
（3）设备端线性存储器的使用说明
虽然CUDA的显存分配函数包括1D，2D和3D的形式，但均不支持多下标访问。对于1D线性空间采用cudaMalloc进行分配和cudaMemcpy进行数据拷贝，其使用方式与分页内存的方式基本一致，当然添加了数据拷贝的方向的控制。还有个需要注意的地方主机和设备间数据交换会自动同步，而设备与设备间不会，需要使用cudaThreadSynchronize()。
但对于2D和3D则不同，以2D为例，分配的函数为cudaMallocPitch，由于它不支持双下标寻址也不支持二级指针，其实就是cudaMalloc的对齐形式，但数据访问方式有大的改变，须采用标准访问形式，即：
T* pElement = (T*)((char*)BaseAddress + Row * pitch) + Column;
注意指针BaseAddress仍为一级指针，本人测试过若将其声明为二级指针，按道理采用如下访问方式：
T* pElement = (T*)((char*)*BaseAddress + Row * pitch) + Column
此方式在模拟条件下能得到正确结果，但实际设备上无法得到正确输出，这也表征了CUDA的线性存储器本质上与内存的不同。
cudaMemcpy2D是用于2D线性存储器的数据拷贝，函数原型为：
cudaMemcpy2D( void* dst，size_t dpitch，const void* src，size_t spitch，size_t width，size_t height，enum cudaMemcpyKind kind )
这里需要特别注意width与pitch的区别，width是实际需要拷贝的数据宽度而pitch是2D线性存储空间分配时对齐的行宽，而当数据传递发生在设备与主机之间时，主机端pitch==width。
综上我们可以看到，CUDA下对二维线性空间的访问是不提供多下标支持的，访问时依然是通过计算偏移量得到，不同的地方在于使用pitch对齐后非常利于实现coalesce访问。
（4）设备端CUDA array的使用方式
CUDA array是专为纹理获取使用的，是不能通过其它方式进行访问的，其分配函数包括cudaMallocArray和cudaMalloc3DArray，这里以2D纹理使用为例，给出一个使用范例。
二维纹理声明：texture tex;
纹理通道声明：cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc ();
CUDA Array的分配与数据拷贝：
cudaMallocArray( &d_array, &channelDesc, width, length );
cudaMemcpyToArray( d_array, 0, 0, d_image, h_size*sizeof(int), cudaMemcpyDeviceToDevice);
纹理参数设置：
tex.addressMode[0] = cudaAddressModeClamp;
tex.addressMode[1] = cudaAddressModeClamp;
tex.filterMode = cudaFilterModePoint;
tex.normalized = false;
纹理绑定：cudaBindTextureToArray(tex, d_array, channelDesc);
二维纹理拾取（未归一化）：tex2D(tex, x, y);
解除纹理绑定：cudaUnbindTexture(tex);