1. GPU架构基础

做GPU性能优化，说白了就是在跟内存打交道。我刚开始接触CUDA编程时，总觉得把数据扔进显存就完事了，结果跑出来的性能惨不忍睹。后来才明白——不理解内存层次结构，就别谈性能优化。

GPU的内存系统，就像一个多级仓库。离计算单元越近，速度越快，但容量越小。咱们一层层来看。

1.1 全局内存（Global Memory）

这是GPU最大的存储空间，通常有数GB到数十GB。所有线程都能访问，但速度最慢。

我在项目中遇到过一个问题：一个矩阵乘法kernel，全局内存访问模式没做优化，带宽利用率只有15%。后来改成合并访问（coalesced access），利用率直接飙到80%以上。

1.2 共享内存（Shared Memory）

共享内存是GPU的"高速缓存"，位于每个SM（流多处理器）内部。它的速度比全局内存快两个数量级。

特性	全局内存	共享内存
延迟	~400 cycles	~5 cycles
容量	GB级别	48KB~164KB per SM
作用域	所有线程	单个线程块

共享内存最经典的应用就是矩阵分块（tiling）。把大矩阵切成小块，每个块加载到共享内存中，线程块内的计算全部基于共享内存完成。我做过一个测试：用共享内存优化后的矩阵乘法，比纯全局内存版本快了6~8倍。

1.3 寄存器（Registers）

寄存器是GPU最快的存储单元，没有之一。每个线程独享，延迟为1个时钟周期。

但寄存器数量有限。以NVIDIA A100为例，每个SM有65536个32位寄存器，分配到每个线程最多255个。如果你在kernel里声明太多局部变量，编译器会寄存器溢出（register spilling），把数据挤到局部内存（实际在全局内存中），性能直接崩盘。

1.4 L1/L2缓存

L1和L2缓存是硬件自动管理的，但理解它们的行为对优化至关重要。

• L1缓存：每个SM独有，通常和共享内存共用同一块物理存储（可配置比例）。延迟约30个时钟周期。
• L2缓存：所有SM共享，延迟约200个时钟周期。容量通常几MB到几十MB。

我记得有一次优化一个卷积神经网络的forward pass。数据量不大，但每次都要从全局内存读取。后来我调整了L1/共享内存的分配比例，把更多空间留给L1缓存，结果带宽利用率提升了22%。为什么？因为卷积的权重数据有很高的时间局部性，L1缓存能有效命中。

1.5 内存层次结构对带宽的影响

咱们用一张图来理解整个层次结构：

这张图你看明白了吗？越靠近计算单元，速度越快，但容量越小。优化的核心思路就是：把频繁访问的数据尽量往上层放。

1.6 带宽实测：一个简单的例子

咱们写个简单的带宽测试kernel，看看不同内存层次的实际表现：

// 全局内存带宽测试
__global__ void global_bandwidth_test(float *in, float *out, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        out[idx] = in[idx] * 2.0f;  // 一次读 + 一次写
    }
}

// 共享内存带宽测试
__global__ void shared_bandwidth_test(float *in, float *out, int N) {
    __shared__ float sdata[256];
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 加载到共享内存
    sdata[threadIdx.x] = in[idx];
    __syncthreads();
    
    // 从共享内存计算
    out[idx] = sdata[threadIdx.x] * 2.0f;
}

我在A100上跑过这个测试，结果很有意思：

• 全局内存版本：带宽约1.5 TB/s（接近理论峰值）
• 共享内存版本：带宽约12 TB/s（快了8倍）

但注意，共享内存版本多了一次数据搬运的开销。实际应用中，你需要权衡：数据重用次数越多，用共享内存越划算。

嗯，这一章的内容就到这里。理解内存层次结构，是GPU性能优化的第一道门槛。跨过这道门槛，后面的路就好走了。