CUDA编程中内存管理机制

GPU设备端存储器的主要分类和特点：

大小：

   全局(Global)和纹理(Texture)内存：大小受RAM大小的限制。

   本地(local)内存：每个线程限制在16KB

   共享内存：最大16kB

   常量内存：总共64KB

   每个SM共有8192或者16384个32位寄存器

速度：

   Global,local,texture << constant << shared,register

数据对齐：

   设备可以在一次操作中从全局内存读取4-byte，8-byte或者16-byte内容到寄存器中，读取不对齐的8-byte或者16-byte内容的可能产生错误的结果。

如何利用合并访问提高访存效率：

   1、使用数组结构体（structure of arrays：SOA）代替结构体数组（array of structures：AOS）：

2、使用共享内存来实现合并访问。

 

内存衬底（memory padding）：

通常的访问模式：二维数组

   当一个索引为(tx,ty)的线程去访问一个宽度为N且基地址为BaseAddress的二维数组时，使用的是下面的地址：BaseAddress + N*ty + tx。在这种情况下，我们如何来保证合并访问呢：

   blockDim.x = 16x 并且 N= 16x。

我们可以控制blockDim.x，但是数组宽度并不总是16x。内存衬底就是创建一个宽度为16x的数组，然后将未使用部分填充0。这里介绍一个概念：数组A的主要尺寸（leading dimension）——pitch，简称Ida。因为c/c++是行主导的，所以主要尺寸为行宽（即一行里面的元素个数）。cuda提供了相应的API，cudaMallocPitch()来分配2D数组。类似的函数同样存在于3D的情况。