谈谈显存的使用分析(二)

二、内存与显存的数据传输类型与使用说明
(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);
展开阅读全文

没有更多推荐了,返回首页