cuda内存的相关理解整理
1.关于各存储的理解
寄存器:寄存器对每个线程是私有的,寄存器是SM上的稀有资源,所以当每个线程使用的寄存器越少,SM能组织的线程就越多,效率也就越高。当寄存器存不下线程的私有变量时,就会使用local memory也就是本地内存,这对效率有巨大的影响,需要尽量避免。
共享内存:共享内存是片上内存,所以速度较快,其的生命周期与块Block是一致的,当block创建,共享内存被分配,block结束,共享内存被回收。所以同样,如果每个线程占用的共享内存过多,会导致组织的线程数量过少,效率大大降低。
常量内存:主机端可以修改但是设备端不可以进行修改
全局内存:最大的内存空间,同时也是最慢的,要求一次读取指定大小(32或者128字节)整数倍的内存,所以必须考虑内存对齐。
这里比较重要的一点是寄存器和共享内存是片上的,所以速度较快,虽然本地内存是针对线程而言的,但是实质上和全局内存是一样的,访问速度很慢
2.零拷贝内存->统一虚拟寻址->统一内存寻址:
零拷贝内存:一般情况下,计算是现将数据从Host端拷贝数据到Device端,然后再device端计算,再将数据存储到主机里,host端和device端的内存各有各的组织形式,而零拷贝内存使得主机和设备都可以访问同一片内存,这片内存存储在主机主存里面,但是可以被设备访问
这里分配了三块零拷贝内存,A和B用来存放输入,C用来存放输出,使用的是主机端的代码进行的计算,然后使用cudaHostGetDevicePointer()函数获取设备到零拷贝内存的指针,比较结果。如下图,两次结果一致,程序也没有出错
统一虚拟寻址:对于零拷贝内存,设备访问零拷贝内存虽然不需要将内存拷贝到设备内存中,但是仍然需要使用新的指针和主机指针指向同一块内存,所以会有很多指针,逻辑会很混乱,统一虚拟寻址使得device端和host端使用统一的指针,精简逻辑的同时也提高了编程的效率。
统一内存寻址:使用cudaMallocManaged()函数分配内存,创建了托管内存池,由底层系统自动分配统一内存,在统一的空间中自动进行设备和主机的数据传输,对用户来说是透明的。
3. 对齐与合并内存访问
核函数运行时需要从全局内存读取数据,有两种粒度:32字节(使用一级缓存)和128字节(不使用一级缓存),将一次内存请求,也就是从内核函数发起请求,到硬件响应返回数据这个过程称为一个内存事务。
当内存事务的首地址是缓存粒度的偶数倍时,被称为对齐内存访问
当访问的内存都在一个内存块里,就会出现合并内存访问
对齐合并内存访问是最高效的。
上图所需的数据存放在以地址128开头的连续128字节的内存块中,所以只需要一次内存事务就可以读取。
如图,读取数据的起始位置不是粒度的倍数,就图中的例子需要使用三次事务读取所需数据,效率大大降低。