CUDA编程之存储模型

CUDA编程之存储模型

CUDA存储模型概述

一般来说，应用程序不会在任何时间点访问任意数据或运行任意代码。程序获取资源是有规律的，也就是计算机体系结构经常提到的局部原则：时间局部性和空间局部性。

• 时间局部性：如果在某时刻访问了某数据，很小可能在段时间内还会访问该数据。
• 空间局部性：如果某时刻访问了某数据，则下一时刻很可能访问与之相邻的数据。
• 总体：如果速度越快，容量越小。

分级存储器体系

以下两张图是GPU在硬件层次与软件层次上对应的存储模型：

内存模型的硬件结构

内存模型的软件结构

全局存储器(Device Memory)

• global Memory是空间最大，latency最高
• 显存
• 慢：400-600 Clocks

纹理缓存(Texture Cache)

• texture Memory实际上也是global Memory在一块，但是他有自己专有的只读cache。
• 加速随机访问
• 具有插值功能
• texture Memory是针对2D空间局部性的优化策略，所以thread要获取2D数据就可以使用texture Memory来达到很高的性能。
• 主要用于图形图像的存储

常量缓存(Constant Cache)

• 加速不变量访问，const restrict*
• 属于全局内存，大小64KB
• 线程请求同一个数据时很快，请求不同的数据时性能下降
• 在运行中不变，所有constant变量的值必须在kernel启动之前从host设置

局部存储器(Local Memory)

• 在local memory中的变量本质上跟global memory在同一块存储区。
• local memory有很高的latency和较低的bandwidth。
• 在CC2.0以上，GPU针对local memory会有L1（per-SM）和L2（per-device）两级cache。

共享存储器(Shared Memory)

• SM中的内存空间(On Chip Memory)
• 作用域是线程块
• 快：30-40 Clocks
• 16/32/48 KB
• 获取shared memory的数据前必须先用__syncthreads()同步

寄存器

• 最快，不同的计算能力数量不同：在Fermi每个thread最多63个registers。
• 如果kernel使用的register超过硬件限制，这部分会使用local memory来代替register，即所谓的register spilling。
• 寄存器是每个thread的私有变量，一旦thread执行结束，寄存器变量就会失效。

变量类型限定符

用来限定变量存储什么位置上。

• __device__表明声明的数据存放在显存中，所有的线程都可以访问，而且主机也可以通过运行时库访问。
• __shared__表示数据存放在共享存储器中，只有所在的 块内的线程可以访问，其它块内的线程不能访问。
• __constant__表明数据存放在常量存储器中，可以被所 有的线程访问，也可以被主机通过运行时库访问。
• texture表明被其绑定的数据可以被纹理缓存加速读取。
• 如果变量没有限定符，那表示它存放在寄存器或者本地存储器中，只归线程所有，其它线程不可见。

总结

变量存储结构总结

参考

CUDA编程指南5.0 [【CUDA】学习记录（7）- Global Memory] https://www.jianshu.com/p/3d4c9cc3a777