📚 课程目录第1章课程导入

GPU内存层次结构概览：从全局内存到寄存器文件

GPU性能优化的核心瓶颈往往不在计算单元，而在内存系统。现代GPU拥有数万个并发线程，它们对数据的访问模式直接决定了内存带宽的利用效率。理解GPU的内存层次结构，是进行任何访存优化的前提。

1.1 GPU内存层次结构全景

以NVIDIA Ampere架构为例，GPU内存系统从物理距离和访问速度上可分为以下几个层级：

寄存器文件（Register File）：每个线程私有的最快存储，容量极小（每线程最多255个32位寄存器），但访问延迟为0周期（若发生bank冲突则产生延迟）。
共享内存（Shared Memory）：每个线程块（Block）内共享，容量通常为48KB~164KB（可配置），延迟约20~30周期，带宽极高。
L1缓存/纹理缓存：与共享内存共用物理SRAM，用于缓存全局内存访问，延迟约30~40周期。
L2缓存：GPU全局共享，容量数MB，延迟约200~300周期，是所有SM（流式多处理器）的汇聚点。
全局内存（Global Memory）/显存（VRAM）：容量最大（数GB~数十GB），但延迟高达400~800周期，带宽受限于显存位宽与频率。
常量内存与纹理内存：具有专用缓存路径，适合广播访问和空间局部性访问。

  关键认知：GPU的访存延迟远高于计算延迟。一个全局内存访问（约500周期）期间，GPU理论上可执行数百条浮点运算指令。因此，隐藏延迟和减少访存次数是优化的两大核心方向。

1.2 优化目标：三个核心指标

GPU内存优化的最终目标可以量化为以下三个指标，它们相互关联、相互制约：

指标	定义	理想值	优化手段示例
内存带宽利用率	实际有效数据传输量 / 理论峰值带宽	> 80%	合并访问、向量化加载（float4）、避免非对齐访问
访存延迟隐藏率	被计算/其他访存掩盖的延迟比例	> 90%	提高占用率、增加每线程独立访存批次、异步预取
片上数据复用率	数据从片上（共享内存/寄存器）读取次数 / 从全局内存读取次数	> 4x	分块（Tiling）、寄存器缓存、共享内存暂存

实践提示：使用NVIDIA Nsight Compute的Memory Workload Analysis视图，可以快速查看当前内核的DRAM Throughput、L2 Cache Hit Rate、Shared Memory Bank Conflicts等关键指标，定位瓶颈属于带宽受限（Bandwidth-bound）还是延迟受限（Latency-bound）。

1.3 学习路径与章节地图

本课程将按照“从微观到宏观、从原理到实践”的路径展开，共30章，分为五大模块：

基础篇（第1~6章）：内存层次结构、合并访问原理、Bank冲突机制、缓存行为分析。
进阶篇（第7~14章）：共享内存分块策略、寄存器缓存、向量化访存、只读缓存与纹理路径。
高级篇（第15~22章）：异步内存拷贝、流水线重叠、持久化线程、内存池与分配器优化。
场景篇（第23~27章）：矩阵乘法、卷积、归约、稀疏访存、图神经网络等典型负载的专项优化。
工具与实战篇（第28~30章）：Nsight Compute深度分析、性能剖析方法论、真实项目优化复盘。

每一章都会围绕一个核心主题，给出可验证的代码示例和性能数据。建议读者在阅读时配合GPU硬件（或云GPU实例）动手实验，因为内存优化的效果高度依赖具体硬件架构。

1.4 常见误区与初始认知纠偏

误区1：所有访存慢的问题都靠加缓存解决。
实际上，GPU的L1/L2缓存命中率提升空间有限（通常L2命中率在30%~60%），且缓存未命中时的惩罚极高。更可靠的方式是从算法层面减少全局内存访问次数，例如通过分块复用数据。

误区2：共享内存越大越好。
共享内存与L1缓存共用SM上的SRAM资源。将共享内存配置调大（如164KB）会压缩L1缓存容量，可能降低全局内存的缓存命中率。需要根据内核访存特征动态权衡。

误区3：线程数越多延迟隐藏越好。
当线程数超过一定阈值后，寄存器溢出（Spill）会导致局部变量被存储到全局内存，反而增加访存压力。每个SM的线程数应控制在寄存器资源允许的范围内。

1.5 一个直观的带宽对比实验

为了建立对内存层次性能差异的直觉，我们可以在任意CUDA设备上运行以下微基准测试，测量不同内存类型的有效带宽：

// 测量全局内存、共享内存、寄存器文件的带宽差异
__global__ void bandwidth_test(float* g_data, float* g_result) {
    __shared__ float s_data[1024];
    int tid = threadIdx.x;
    int gid = blockIdx.x * blockDim.x + tid;
    
    // 1. 全局内存读取
    float g_val = g_data[gid];
    
    // 2. 共享内存写入+读取
    s_data[tid] = g_val;
    __syncthreads();
    float s_val = s_data[(tid + 1) % 1024];
    
    // 3. 寄存器操作
    float r_val = g_val * 2.0f + s_val;
    
    g_result[gid] = r_val;
}

在A100上运行该内核，你会观察到：全局内存有效带宽约1.5~1.8TB/s（理论峰值2TB/s的75%~90%），共享内存带宽可达10~15TB/s，而寄存器操作带宽更高。这个数量级差异正是所有优化技巧的出发点——让数据尽量在片上流动，减少与片外显存的交互。