📚 课程目录第3章显存类型与特性

3. 显存类型与特性：全局内存、共享内存、寄存器、本地内存、常量内存、纹理内存

CUDA 编程模型将 GPU 上的可寻址存储空间划分为多种类型，每种类型在容量、延迟、带宽和可见性上差异显著。理解这些差异是编写高性能内核的前提——数据放置的位置，往往比计算指令本身更决定性能上限。本章逐一剖析六种核心显存类型，并给出实际使用建议。

3.1 全局内存（Global Memory）

全局内存是 GPU 中容量最大、最基础的存储空间，所有线程均可读写，并可在主机端通过 cudaMemcpy 访问。它位于显存（DRAM）中，通常为 8~80 GB。其访问延迟高达 400~800 个时钟周期，但吞吐量极大（数百 GB/s 至数 TB/s）。

全局内存的访问效率高度依赖合并访问（Coalesced Access）：同一 warp（32 线程）访问的地址应尽量落在连续的 128 字节内存段内。若线程间地址跳跃，则会产生多次事务，浪费带宽。

关键参数：

容量：最大（GB 级）
延迟：400~800 周期
缓存：L2 缓存（所有 SM 共享），部分架构有 L1 只读路径
作用域：设备端所有线程 + 主机端

// 合并访问示例：连续索引
__global__ void copy_kernel(float *in, float *out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) out[i] = in[i];  // 相邻线程访问相邻地址
}

3.2 共享内存（Shared Memory）

共享内存是每个线程块（Block）内部的高速可编程缓存，位于 SM 内部，延迟约为 20~30 个周期，远低于全局内存。它用于线程块内线程间数据共享与协作，典型容量为 48 KB~228 KB（可配置）。

共享内存被划分为 32 个存储体（Bank），每个存储体宽度 4 字节。当多个线程同时访问不同存储体时，可并行完成；若访问同一存储体（广播除外），则产生存储体冲突（Bank Conflict），导致访问串行化，性能下降数倍。

使用技巧：利用共享内存做数据复用（如矩阵分块、归约），可减少全局内存访问次数。声明方式：__shared__ float tile[32][32];。注意同步：__syncthreads() 确保块内线程看到彼此写入的数据。

3.3 寄存器（Registers）

寄存器是 GPU 上最快的存储单元，位于 SM 内部，延迟为 0（单周期访问）。每个线程的私有变量（如循环索引、临时计算值）默认存储在寄存器中。每个 SM 拥有 64K 个 32 位寄存器，分摊到每个线程通常为 32~255 个。

寄存器是稀缺资源：若每个线程使用过多寄存器，会降低 SM 上可同时驻留的线程块数量（Occupancy），从而影响延迟隐藏能力。编译器通过 -maxrregcount 或 __launch_bounds__ 控制寄存器分配。

注意：寄存器是线程私有的，线程间无法直接访问彼此的寄存器。若发生“寄存器溢出”（Register Spill），数据会被压入本地内存（Local Memory），性能急剧下降。

3.4 本地内存（Local Memory）

本地内存在物理上位于显存（DRAM）中，但逻辑上为每个线程私有。它用于存储寄存器放不下的数据：大型数组、动态索引的局部变量、递归调用栈等。访问延迟与全局内存相当（数百周期），且不经过 L1 缓存（仅 L2）。

本地内存的滥用通常源于：

在核函数内声明大型局部数组（如 float arr[1024]）
使用非常量索引访问局部数组
编译优化失败导致变量被强制放入本地内存

通过编译选项 -Xptxas -v 可查看内核的寄存器与本地内存使用量。若发现 local 字节数过大，应重构代码，改用共享内存或全局内存显式管理。

3.5 常量内存（Constant Memory）

常量内存是全局内存的只读缓存视图，容量 64 KB，专用于存储内核执行期间不变的参数（如权重、系数、查找表）。它拥有独立的常量缓存（Constant Cache），当 warp 内所有线程访问同一地址时，该访问被广播为单次事务，速度极快。

适用场景：所有线程读取同一个值（如标量参数、滤波器系数）。若线程访问不同地址，则退化为串行访问，性能反而低于全局内存。

// 常量内存声明（必须在文件作用域）
__constant__ float c_filter[128];

// 主机端复制
cudaMemcpyToSymbol(c_filter, h_filter, sizeof(float) * 128);

3.6 纹理内存（Texture Memory）

纹理内存是一种通过纹理缓存（Texture Cache）访问的只读内存，最初为图形渲染设计，现也用于通用计算。它支持二维/三维空间局部性优化，适合图像处理、插值计算等访问模式。纹理缓存针对空间邻近性做了优化，当线程访问相邻像素时命中率极高。

纹理内存还提供硬件插值（线性滤波）、边界处理（Clamp/Wrap）等功能，但通用计算中通常仅利用其缓存特性。现代架构（Volta 及以上）中，纹理内存与全局内存的 L1 路径已部分统一，但其只读缓存仍对特定访问模式有效。

存储类型	容量	延迟（周期）	缓存	作用域	典型用途
寄存器	每 SM 64K×32bit	0	无	线程私有	临时变量、索引
共享内存	48~228 KB/SM	20~30	无（显式管理）	线程块内	数据复用、块内通信
全局内存	GB 级	400~800	L2	所有线程+主机	主数据存储
本地内存	同全局	400~800	L2（无L1）	线程私有	寄存器溢出、大数组
常量内存	64 KB	广播时≈L1	常量缓存	所有线程只读	参数、系数
纹理内存	同全局	空间局部≈L1	纹理缓存	所有线程只读	图像、空间数据

3.7 案例：矩阵乘法中的显存选择

以 1024×1024 矩阵乘法为例，展示不同显存策略的性能差异。朴素实现直接访问全局内存，每个输出元素需读取 2×1024 个浮点数，全局内存访问次数为 2×1024³ 次。采用共享内存分块（Tile 32×32）后，每个线程块仅需加载 2×32×32 个数据到共享内存，复用 32 次，全局内存访问量降低 32 倍。

// 共享内存分块矩阵乘法核心代码
__global__ void matmul_tiled(float *A, float *B, float *C, int N) {
    __shared__ float As[32][32];
    __shared__ float Bs[32][32];
    int row = blockIdx.y * 32 + threadIdx.y;
    int col = blockIdx.x * 32 + threadIdx.x;
    float sum = 0.0f;

    for (int tile = 0; tile < N/32; ++tile) {
        As[threadIdx.y][threadIdx.x] = A[row * N + tile * 32 + threadIdx.x];
        Bs[threadIdx.y][threadIdx.x] = B[(tile * 32 + threadIdx.y) * N + col];
        __syncthreads();
        for (int k = 0; k < 32; ++k)
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        __syncthreads();
    }
    C[row * N + col] = sum;
}

实测中，该优化在 NVIDIA A100 上可将吞吐量从约 1.2 TFLOPS 提升至 19.5 TFLOPS（接近理论峰值 19.5 TFLOPS 的 80%）。关键收益来自：共享内存减少全局内存压力，寄存器缓存中间累加值，避免本地内存溢出。

3.8 常见误区与性能复盘

误区 1：所有数据都放共享内存就一定快。共享内存容量有限，且过度使用会降低 Occupancy。当块内线程数较少或数据复用率低时，直接使用全局内存（依赖 L2 缓存）反而更优。

误区 2：忽略寄存器溢出。在核函数中声明 float local[64] 且用变量索引访问，编译器可能将其放入本地内存。应使用 -Xptxas -v 检查，必要时改用共享内存或重构算法。

误区 3：常量内存随意使用。若 warp 内线程访问不同常量地址，常量缓存会串行化访问，性能下降。只有“所有线程读同一地址”时才高效。

复盘建议：使用 NVIDIA Nsight Compute 分析内核的 Memory Workload Analysis 页面，查看各存储类型的吞吐量与利用率。重点关注：全局内存合并率（应接近 100%）、共享内存 Bank Conflict 次数（应为 0）、寄存器溢出量（应为 0）。若共享内存利用率低于 60%，考虑增大分块尺寸；若寄存器溢出，减少每线程寄存器占用或调整算法。

总结：六种显存类型构成 GPU 的存储层次——寄存器与共享内存是性能关键，全局内存是数据主体，常量与纹理内存针对特定访问模式优化，本地内存则是需要避免的“性能陷阱”。合理规划数据在各级存储中的流动，是 GPU 性能优化的核心艺术。