57 lines
2.7 KiB
Markdown
57 lines
2.7 KiB
Markdown
# 全局内存的合理使用
|
||
|
||
在各种设备内存中，全局内存具有最低的访问速度，往往是一个 CUDA 程序的性能瓶颈。
|
||
|
||
------
|
||
|
||
## 全局内存的合并与非合并访问
|
||
|
||
对全局内存的访问将触发内存事务，即数据传输。
|
||
在启用了 L1 缓存的情况下，对全局内存的读取将首先尝试经过 L1 缓存；如果未命中，
|
||
则尝试经过 L2 缓存；如果再次未命中，则直接从 DRAM 读取。
|
||
|
||
一次 **数据传输处理** 的数据量在默认情况下是 32 字节。
|
||
一次数据传输中，从全局内存转移到 L2 缓存的一片内存的首地址一定是 32 的整数倍。
|
||
也就是说，一次数据传输只能从全局内存读取地址为 0-31 字节、32-63 字节等片段的数据。
|
||
|
||
**合并度**，即线程束请求的字节数与由此导致的所有内存事务中所传输的字节数之比。
|
||
如果所有数据传输处理的数据都是线程束所需要的，则合并度为 100%，即 **合并访问**；
|
||
否则，即为 **非合并访问**。
|
||
|
||
以仅使用 L2 缓存的情况为例，一次数据传输指的就是将 32 字节数据从全局内存（DRAM）
|
||
通过 32 字节的 L2 缓存片段（cache sector）传输到 SM。
|
||
考虑一个线程束访问单精度浮点数类型的全局内存变量的场景，
|
||
一个单精度浮点数占有 4 个字节，故一次访问需要 32*4 个字节的数据。在理想情况下，
|
||
即合并度为 100% 时，将仅触发 128/32=4 次调用 L2 缓存的数据传输。
|
||
如果线程束请求的全局内存地址刚好为 0-127 字节或 128-255 字节，就能与 4 次数据
|
||
传输所处理的数据完全吻合，这种情况下就是合并访问。
|
||
|
||
64 位系统中基本数据类型的内存长度（字节）：
|
||
|
||
int size: 4
|
||
short size: 2
|
||
float size: 4
|
||
double size: 8
|
||
char size: 1
|
||
bool size: 1
|
||
long size: 4
|
||
int pointer size: 8
|
||
float pointer size: 8
|
||
double pointer size: 8
|
||
char pointer size: 8
|
||
|
||
------
|
||
|
||
## 矩阵转置
|
||
|
||
在核函数中，如果读取操作是非合并访问，则可以采用 *只读数据缓存技术*，通过加载函数
|
||
`__ldg()` 读取全局内存，从而对数据的读取进行缓存、缓解非合并访问的影响。
|
||
|
||
从帕斯卡架构开始，编译器会自动判断并调用 `__ldg()` 函数提升性能；对于开普勒架构、
|
||
麦克斯韦架构，默认情况下不会使用 `__ldg()` 函数，需要手动配置。
|
||
|
||
对于核函数中全局内存的写入，则没有类似函数可用。所以若不能满足读取和写入都是合并的，
|
||
一般应该尽量做到合并写入。
|
||
|
||
------
|