GPU内存体系及其优化指南
GPU内存体系及其优化指南
要在冯·诺依曼架构的硬件中实现高性能计算，最重要的两点就是：访存和计算**。这两点分别对应着 IO bound 和 compute bound ，硬件系统的内存体系深刻影响着这两点。因此要实现软件层面的高性能计算，必须要对内存体系有深刻的理解。本篇主要讨论GPU的内存体系，并在此基础上进行CUDA的编程实践。
一、GPU的内存体系
内存的访问和管理编程语言的重要组成部分，也是实现高性能计算的重要环节。CUDA 内存模型结合了主机和设备的内存系统，具有完整的层次结构，并可以显式地进行控制和优化。
1.1 各级内存及其特点
CUDA 内存模型的层次结构
全局内存(global memory)
全局内存是GPU中最大、延迟最高、最长使用的内存，通常说的“显存”中的大部分都是全局内存。全局内存的声明可以在任何SM设备上被访问到，并且贯穿应用程序的整个生命周期。
全局内存的主要角色是为核函数提供数据，并在主机与设备及设备与设备之间传递数据。可以用
cudaMemcpy
函数将主机的数据复制到全局内存，或者反过来。如将中 M 字节的数据从主机复制到设备，操作如下：
cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);
全局内存变量可以被静态声明和动态声明， 如 静态全局内存变量由以下方式在任何函数外部定义 ：
__device__ T x; // 单个变量
__device__ T y[N]; // 固定长度的数组
后续将会重点研究如何优化全局内存访问，以及如何提高全局内存的数据吞吐率。
常量内存(constant memory)
常量内存是指存储在片下存储的设备内存上，但是通过特殊的常量内存缓存（constant cache）进行缓存读取，常量内存为只读内存。常量内存数量有限，一共仅有 64 KB，由于有缓存，常量内存的访问速度比全局内存高，但得到高访问速度的前提是一个线程束中的线程(一个线程块中相邻的 32 个线程)要读取相同的常量内存数据。
一个使用常量内存的方法是在核函数外面用
__constant__
定义变量，并用 API 函数
cudaMemcpyToSymbol
将数据从主机端复制到设备的常量内存后 供核函数使用。
纹理内存(texture memory)和表面内存(surface memory)
纹理内存和表面内存类似于常量内存，也是一 种具有缓存的全局内存，有相同的可见范围和生命周期，而且一般仅可读(表面内存也可 写)。不同的是，纹理内存和表面内存容量更大，而且使用方式和常量内存也不一样。
寄存器（register）
寄存器是线程能独立访问的资源，它所在的位置与局部内存不一样，是在片上（on chip）的存储，用来存储一些线程的暂存数据。寄存器的速度是访问中最快的，但是它的容量较小。
在核函数中定义的不加任何限定符的变量一般来说就存放于寄存器(register)中。各种内建变量，如 gridDim、blockDim、blockIdx、 threadIdx 及 warpSize 都保存在特殊的寄存器中，以便高效访问。在上期求和的例子中：
const int n = blockDim.x * blockIdx.x + threadIdx.x;
c[n] = a[n] + b[n];
中的 n 就是一个寄存器变量。寄存器变量仅仅被一个线程可见。也就是说，每一个线程都有一个变量 n 的副本。虽然在核函数的代码中用了这同一个变量名，但是不同的线程中该寄存器变量的值是可以不 同的。每个线程都只能对它的副本进行读写。寄存器的生命周期也与所属线程的生命周期 一致，从定义它开始，到线程消失时结束。
局部内存(local memory)
局部内存和寄存器几乎一 样，核函数中定义的不加任何限定符的变量有可能在寄存器中，也有可能在局部内存中。寄存器中放不下的变量，以及索引值不能在编译时就确定的数组，都有可能放在局部内存中。
虽然局部内存在用法上类似于寄存器，但从硬件来看，局部内存只是全局内存的一部 分。所以，局部内存的延迟也很高。每个线程最多能使用高达 512 KB 的局部内存，但使用 过多会降低程序的性能。
共享内存(shared memory)
共享内存和寄存器类似，存在于芯片 上，具有仅次于寄存器的读写速度，数量也有限。一个使用共享内存的变量可以
__shared__
修饰符来定义。
不同于寄存器的是，共享内存对整个线程块可见，其生命周期也与整个线程块一致。也 就是说，每个线程块拥有一个共享内存变量的副本。共享内存变量的值在不同的线程块中 可以不同。一个线程块中的所有线程都可以访问该线程块的共享内存变量副本，但是不能 访问其他线程块的共享内存变量副本。共享内存的主要作用是减少对全局内存的访问，或 者改善对全局内存的访问模式。
以上内存的主要特点如下表所示：
L1和L2 缓存
每个 SM 都有一个 L1 缓存，所有 SM 共享一个 L2 缓存。L1 和 L2 缓存都被用来存储局部内存和全局内存中的数据，也包括寄存器中溢出的部分，以减少延时。
从物理结构上来说，在最新的GPU架构中，L1 缓存、纹理缓存及共享内存三者是统一的。但从编程的角度来看，共享 内存是可编程的缓存(共享内存的使用完全由用户操控)，而 L1 和 L2 缓存是不可编程的缓存(用户最多能引导编译器做一些选择)。
1.2 SM 构成及典型GPU的对比
一个 GPU 是由多个 SM 构成的。一个 SM 包含如下资源:
一定数量的寄存器。
一定数量的共享内存。
常量内存的缓存。
纹理和表面内存的缓存。
L1缓存。
线程束调度器(warp scheduler) 。
执行核心，包括:
若干整型数