CUDA设备端线性存储器

CUDA学习之内存分配

一、CUDA线性内存分配

1. 1D线性内存分配

cudaMalloc(void **, int);       //在设备端分配内存
cudaMemcpy(void* dest, void* source, int size, enum direction);   //数据拷贝
cudaMemcpyToSymbol   //将数据复制到__constant__变量中，或者__device__变量中
cudaMemcpyFromSynbol   //同上相反
cudaFree()               //内存释放
cudaMemset()           //内存初始化

cudaMalloc
参数1：GPU设备端数据指针
参数2：空间大小，字节为单位
cudaMemcpy
参数1：目标数据地址
参数2：源数据地址
参数3：数据大小
参数4: 方向：cudaMemcpyHostToDevice，主机到设备； cudaMemcpyDeviceToHost，设备到主机

注意：主机和设备间的数据交换会自动同步，而设备与设备却不会，需要使用cudaThreadSynchronize()，强制等待所有在此句之前启动的runtime任务完成。此后再进行下一个任务。

2. 2D线性内存分配
2.1 分配

cudaMallocPitch(void** devPtr, size_t* pitch, size_t widthInBytes, size_t height); //在线性内存中分配二维数组，width 的单位是字节，height单位是数据类型

参数1：设备端数据指针
参数2：一行数据的真实空间大小（字节）【此参数是获取返回值】，实际大小大于需要分配的大小
参数3：每行需要分配的空间大小
参数4： 行数
c语言申请2维内存时，一般是连续存放的。a[y][x]存放在第ywidthofxsizeof(元素)+x*sizeof(元素)个字节。

注意：cudaMallocPitch函数分配的内存中，数组的每一行的每一个元素的开始地址都是保证对齐的。为保证数组的每一行的第一个元素的开始地址对齐，cudaMallocPitch在分配内存时，每行会多分配一些字节，以保证widthofx*sizeof(元素)+多分配的字节是256的倍数(对齐)。这样，上面的ywidthofxsizeof(元素)+xsizeof(元素)来计算a[y][x]的地址就不正确了。而应该是y[widthofxsizeof(元素)+多分配的字节]+xsizeof(元素)。而函数中返回的pitch的值就是widthofx*sizeof(元素)+多分配的字节。说明：widthInBytes作为输入参数，应该是widthofx*sizeof(元素)；这样的话，复制内容时也要作相应的修改。

2.2 访问
如果给定一个T类型数组元素的行和列，可按如下方法计算地址：

T* pElement = (T*)((char*)BaseAddress + Row * pitch) + Column;           //元素访问方式

先计算访问的行的初始地址 （T*）（ (char*)BaseAddress + Row * pitch）
然后访问此行对应元素 （T*）（ (char*)BaseAddress + Row * pitch）+ Column

注意：cudaMallocPitch()以*pitch的形式返回间距，即所分配存储器的宽度，以字节为单位。间距用作存储器分配的一个独立参数，用于在2D数组内计算地址。

2.3 拷贝

cudaMemcpy2D( void* dst，size_t dpitch，const void* src，size_t spitch，size_t width，size_t height，enum cudaMemcpyKind kind )

参数1：目标数据地址

参数2：pitch，分配空间的行宽（字节单位）

参数3：源数据地址

参数4：pitch，分配空间的行宽（字节单位）

参数5：需要拷贝数据的真实行宽（字节单位）

参数6：数据的行数（非字节单位）

参数7：数据拷贝类型

注意：width与pitch的区别，width是实际需要拷贝的数据宽度而pitch是2D线性存储空间分配时对齐的行宽，而当数据传递发生在设备与主机之间时，主机端pitch==width.

3. 3D线性内存分配
敬请期待。。。