CUDA 004 - 内存模型

整体结构

如同 CPU 中有多个内存层级（如图缓存、主存等），GPU 中也有多种类型的内存。理解 GPU 的内存层级对编写高效 GPU 程序至关重要。GPU 的内存层级的构成比 CPU 简单很多，下图简要地描述了 GPU 的内存层级构成：

Register File

每个 partition 有 16,384 个可以存储 4-byte 数据的寄存器，寄存器总量为 64KB，GPU 的寄存器总量相比 CPU 而言要丰富很多。当一个 block 被调度执行时，会给 block 中的所有的 thread 预先分配好寄存器。寄存器实际上就是一块连续的内存，编译器在编译 CUDA 函数时，会将 thread 用到的寄存器分配到一起。这样每个 thread 只需要根据自己的 thread 编号计算出一个偏移量就可以访问属于该 thread 的寄存器。

因为 block 中所有 thread 的寄存器在调度执行 block 时都已经分配好了，因此如果某个 warp 在执行时被阻塞，可以很快速地切换到其他 warp 上执行。和在 CPU 上切换 thread 相比，GPU 上切换 thread 的代价要小很多，它需要保存的现场很少。所以 GPU 可以通过频繁地切换 warp 来降低时延带来的影响。

上图中绿色表示 warp 在执行的时间片，红色表示被阻塞的时间片。虽然每个 warp 都有较长时间的阻塞，但是通过切换 warp 可以让 GPU 充分利用好它的计算资源，避免 idle 状态。

L1 data cache / Shared Memory

L1 data cache 和 shared memory 是每个 SM 私有的，它们共用一个物理内存空间的，如果 shared memory 消耗的容量比较大，那么 L1 data cache 可使用的就比较小。如果没有使用 shared memory，那么 L1 data cache 可以使用全部的内存。

L1 data cache 和 shared memory 的访问延迟大约 20 - 30 个时钟周期，相比访问 global memory 的延迟要小很多。作为对比，CPU 的 L1 cache 访问延迟大约 5 时钟周期。

Shared Memory

访问 shared memory 的时延远远小于访问 global memory 的时延。对一些需要频繁修改的数据，可以将其存放在 shared memory 中，在计算完成后再将结果写入到主存。另外，当 block 中的多个 thread 之间需要共享数据时，可以将数据存放在 shared memory 中。

在 CUDA 核函数中使用 __shared__ 修饰一个变量或数组后，这个变量就会被分配到 shared memory 中，在整个 block 中所有 thread 均可访问。

Global Memory

Global memory 是 GPU 的主存，就像 CPU 的 RAM 一样。普通消费级显卡的 global memory 大小一般为 8GB-16GB，使用 GDDR 作为存储介质。用在数据中心做科学计算的高端 GPU 会使用 HBM（High Bandwidth Memory）作为存储介质，大小从 40GB-96GB 不等。global memory 可以被 host 和所有 thread 访问，可以使用 cudaMalloc 来动态申请，可以使用 cudaMemcpy 来进行 host 和 device 之间的内存拷贝。

对 global memory 的访问是以 32、64、128 字节的块为单位进行的。如果一个 block 内的 thread 们访问的内存地址是连续的，那么可以一起访问 global memory，这样的操作叫做 coalesced 访问。虽然多个 thread 都发起了全局内存的读写操作，但这些读写操作可以合并起来，通过一个 memory transaction 来完成。

Local Memory

有些文档中会提到 local memory，它指的是 thread 私有访问的内存。Local memory 是一个概念性的内存类型，并不存在一块特别的内存叫做 local memory。如果 thread 的局部变量太多，无法完全放入寄存器中，部分局部变量也会存储在 global memory 中，但被称为 local memory。

可能被编译器存放到 local memory 的变量：

• 使用非常量索引访问的数组，比如 arr[i]，这类数组虽然可能很小，但因为编译器无法判断是否存在访问越界的情况，所以这类数组会被放在 local memory 中。
• 如果变量是会占用用较大内存的大型结构体或者数组，编译器可能会将其放置在 local memory。
• 线程内使用了太多局部变量，寄存器不够用，其他变量会被放置在 local memory。

Constant Memory

使用 __constant__ 修饰的变量会被放置在 global memory 中，但在逻辑上被称为 constant memory 中。constant memory 是只读的，在被访问时，会逐步被加载到 SM 中的 constant cache 中。

总结

GPU 的内存结构同样反映了 GPU 适合进行并行计算的特性。GPU 中有大量的寄存器，这可以让 warp 切换变的更加轻量，遇到访问内存这样的耗时比较长的操作，GPU 内部也可以执行 warp 切换，在数据加载时处理器可以转而处理其他已经就位的数据。

访问全局变量时，GPU 会以最小 32 字节的块为单位进行访问。虽然读取速度比起 CPU 访问主存要慢，但这样可以提高数据读取的吞吐量，而且 GPU 中多个线程之间往往具有空间连续性，加载一次数据，就可以服务于多个线程。

GPU 的内存层级相比 CPU 而言要简单不少。多个线程的数据共享被限制在了 block 内，而 shared memory 和 L1 cache 是 SM 私有的，这就避免引入 cache consistency 的问题。

在编写 CUDA 程序的时候，需要能够想到 GPU 内存的特性，这样才能写出高效的 GPU 程序。