第四章：内存管理与优化——全局内存、共享内存、寄存器、常量内存、纹理内存的使用与优化策略

说到GPU性能优化，内存管理绝对是绕不开的核心话题。我见过太多同学，核函数写得花里胡哨，结果一跑性能惨不忍睹。为什么？说白了，大部分时间都浪费在了数据搬运上。GPU的计算单元快得飞起，但内存访问如果跟不上，那就是典型的「高射炮打蚊子」。

这一章，咱们就把GPU的五种内存掰开揉碎了讲清楚。你想想看，搞懂了这些，你的算子性能至少能翻个倍。

4.1 全局内存：最常用，也最慢

全局内存是GPU里容量最大的存储空间，所有线程都能访问。但它的延迟也是最高的——大概400-800个时钟周期。嗯，这个数字你得记住。

访问模式是关键

全局内存有个脾气：它喜欢「合并访问」。什么意思？就是相邻的线程最好访问相邻的地址。我在项目中遇到过，一个同事写了个矩阵转置，性能死活上不去。我一看代码，好家伙，线程访问的步长是矩阵的行宽，完全没对齐。改成列访问后，带宽利用率从20%直接飙到了80%。

// ❌ 糟糕的访问模式：步长访问
float value = data[threadIdx.x * 1024];  // 每个线程跳着读

// ✅ 良好的访问模式：连续访问
float value = data[threadIdx.x];  // 相邻线程读相邻地址

4.2 共享内存：片上高速缓存

共享内存是GPU上的「小仓库」，位于芯片内部，延迟只有全局内存的十分之一左右。每个线程块独享一份共享内存，块内的线程可以互相通信。

容量有限，精打细算

每个SM（流多处理器）的共享内存通常只有几十到一百多KB。你想想看，这点空间得省着用。我建议你养成习惯：每次写代码前先算算共享内存的用量。

// 声明共享内存
__shared__ float cache[256];

// 使用共享内存做归约
int tid = threadIdx.x;
cache[tid] = input[tid];
__syncthreads();  // 同步！确保所有线程都写完了

// 树形归约
for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
    if (tid < stride) {
        cache[tid] += cache[tid + stride];
    }
    __syncthreads();
}

4.3 寄存器：最快但最稀缺

寄存器是GPU里最快的存储，没有之一。每个线程独享，延迟几乎为零。但数量极其有限——每个线程通常只有32-255个寄存器（取决于GPU架构）。

寄存器溢出是大忌

如果每个线程使用的寄存器超过了硬件限制，编译器会把多余的数据「溢出」到全局内存。这一溢出，性能就崩了。我见过一个极端案例：一个核函数用了128个寄存器，结果编译器溢出后，性能下降了10倍。

// 查看寄存器使用量
nvcc --ptxas-options=-v kernel.cu
// 输出类似：registers=32, smem=0, cmem[0]=16

4.4 常量内存：只读且广播

常量内存是全局内存的一种特殊形式，但有两个特点：一是只读，二是当线程束内所有线程访问同一个地址时，硬件会把这个值广播给所有线程，只发一次请求。

适用场景

常量内存最适合存放那些不会改变、且所有线程都需要的参数。比如卷积核的权重、矩阵乘法的系数等。

// 声明常量内存
__constant__ float filter[1024];

// 主机端复制数据
cudaMemcpyToSymbol(filter, host_filter, sizeof(float) * 1024);

// 核函数中直接使用
__global__ void convolution(float* output, float* input) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    float sum = 0.0f;
    for (int i = 0; i < 1024; i++) {
        sum += input[idx + i] * filter[i];  // 所有线程读同一个filter[i]
    }
    output[idx] = sum;
}

4.5 纹理内存：为空间局部性而生

纹理内存最初是为图形学设计的，但用在通用计算里也有奇效。它利用空间局部性，对二维或三维数据的随机访问做了优化。而且它自带硬件插值功能，做图像处理时特别方便。

什么时候用纹理内存？

• 数据访问模式不规则，比如光线追踪、随机采样
• 需要硬件插值，比如图像缩放、特征提取
• 数据是只读的，且访问模式有空间局部性

// 绑定纹理
texture<float, 2, cudaReadModeElementType> tex;
cudaBindTexture2D(0, tex, dev_data, &tex.channelDesc, width, height, pitch);

// 核函数中采样
__global__ void kernel(float* output) {
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    output[y * width + x] = tex2D(tex, x + 0.5f, y + 0.5f);  // 硬件插值
}

4.6 五种内存对比与选择策略

为了方便你快速决策，我整理了一张表。你写代码前可以对照着看看：

内存类型	位置	延迟	容量	作用域	典型用途
全局内存	片外DRAM	高（~400周期）	大（GB级）	所有线程	主数据存储
共享内存	片上SRAM	低（~30周期）	小（KB级）	线程块内	数据复用、线程通信
寄存器	片上	极低（~1周期）	极小（每个线程几十个）	单个线程	局部变量、中间结果
常量内存	片外（带缓存）	低（广播时）	64KB	所有线程	只读参数、系数
纹理内存	片外（带缓存）	中（空间局部性好）	大（GB级）	所有线程	不规则访问、插值

4.7 知识体系总览

下面这张图是我自己画的，把五种内存的关系和优化策略串在了一起。你保存下来，写代码时对照着看：

4.8 实战建议：如何选择内存类型

写代码时，我一般按这个顺序思考：

1. 数据是否只读？ → 是的话，考虑常量内存或纹理内存
2. 数据是否被反复使用？ → 是的话，用共享内存缓存
3. 访问模式是否连续？ → 是的话，直接用全局内存合并访问
4. 数据量是否很小？ → 是的话，尽量塞进寄存器

说白了，GPU优化的本质就是「让数据尽量靠近计算单元」。寄存器最快但最小，全局内存最慢但最大。你要做的，就是在两者之间找到平衡点。