3、GPU内存层次结构：全局内存、共享内存、寄存器、本地内存、常量内存、纹理内存

聊GPU编程，绕不开的就是内存。我经常跟团队里的新人说，搞懂GPU的内存层次，你就掌握了性能优化的钥匙。说白了，GPU计算的核心就是数据搬运——把数据从慢的地方搬到快的地方，算完再搬回去。你想想看，如果数据搬运的路径没选对，计算单元再快也是白搭。

GPU的内存不像CPU那样简单。它是一套复杂的层次结构，每种内存都有自己脾气。我刚开始接触CUDA时，就吃过不少亏。有一次写一个矩阵乘法的kernel，全局内存访问模式没优化好，性能比CPU还慢。后来改成共享内存加合并访问，速度直接翻了20倍。嗯，这就是内存层次结构的魔力。

3.1 全局内存（Global Memory）

全局内存是GPU里容量最大的存储空间。它位于显存中，所有线程都能访问。但代价是——访问延迟非常高，大概400-800个时钟周期。

我在项目中遇到过一个问题：一个图像处理kernel，每次迭代都要从全局内存读一个像素值。结果发现，大部分时间都花在了等待数据上。后来我做了两件事：一是把数据对齐到128字节边界，二是让相邻线程访问相邻地址。这就是所谓的合并访问（Coalesced Access）。

举个例子，假设你要读一个float数组：

// 好的访问模式：合并访问
float val = data[threadIdx.x + blockIdx.x * blockDim.x];

// 坏的访问模式：非合并访问
float val = data[threadIdx.x * stride + offset];

前者相邻线程访问相邻地址，后者跳着访问。性能差距可能达到10倍以上。

3.2 共享内存（Shared Memory）

共享内存是我个人最喜欢的GPU存储资源。它位于芯片内部，延迟只有几个时钟周期。一个线程块内的所有线程都能访问它。你可以把它想象成一个手动管理的高速缓存。

我记得第一次用共享内存优化矩阵转置时，那种感觉真的很爽。原来需要反复读全局内存的数据，现在一次性搬到共享内存里，然后从共享内存里读。性能提升立竿见影。

共享内存有两个关键特性：

• 容量有限：每个SM通常只有48KB或96KB（取决于架构）。别想着把所有数据都塞进去。
• Bank冲突：共享内存被分成32个bank。如果多个线程同时访问同一个bank的不同地址，就会产生冲突，导致串行化。

使用共享内存的典型模式：

__global__ void kernel(float *input, float *output) {
    __shared__ float s_data[256];
    
    // 第一步：从全局内存加载到共享内存
    int tid = threadIdx.x;
    s_data[tid] = input[tid + blockIdx.x * blockDim.x];
    __syncthreads();  // 同步！确保所有线程都加载完毕
    
    // 第二步：在共享内存上计算
    // ... 做一些计算 ...
    
    // 第三步：写回全局内存
    output[tid + blockIdx.x * blockDim.x] = s_data[tid];
}

注意那个 __syncthreads()。我曾经因为忘记加这个同步，导致一个线程读到的数据还是上一轮的结果。调试了整整一个下午才找到问题。嗯，这个坑我替你们踩过了。

3.3 寄存器（Register）

寄存器是GPU里最快的存储资源。没有之一。每个线程都有自己的寄存器，其他线程无法访问。访问延迟几乎为零。

但寄存器数量非常有限。每个SM的寄存器总数是固定的（比如Kepler架构是65536个），分配到每个线程的寄存器数量取决于线程数。如果一个线程用了太多寄存器，编译器会做两件事：

• 寄存器溢出（Register Spilling）：把多余的变量存到本地内存（Local Memory）里。本地内存其实在显存里，速度会慢很多。
• 降低占用率：每个SM能同时运行的线程块数量减少。

你可以通过 launch_bounds 来提示编译器每个块的最大线程数，从而控制寄存器分配：

__global__ __launch_bounds__(256, 4)  // 最多256线程，最少4个块
void myKernel() {
    // ...
}

3.4 本地内存（Local Memory）

本地内存这个名字有点误导。它其实不在芯片内部，而是位于显存中。编译器会自动把寄存器放不下的变量、数组索引访问的变量、以及一些大型结构体放到本地内存里。

本地内存的访问延迟和全局内存差不多。但它的访问模式是每个线程私有的，所以不会产生合并访问的问题。不过，频繁使用本地内存会严重影响性能。

我曾经优化过一个kernel，发现性能始终上不去。用 nvcc --ptxas-options=-v 一看，本地内存用了不少。原来是代码里定义了一个大数组：

// 这个数组会被放到本地内存
float temp[1024];

改成共享内存后，性能提升明显。所以我的建议是：尽量避免在kernel内部定义大数组。如果必须用，考虑用共享内存或者重新设计算法。

3.5 常量内存（Constant Memory）

常量内存是全局内存的一个特殊区域。它有以下几个特点：

• 只读：所有线程只能读，不能写。
• 有缓存：每个SM都有一个常量缓存，大小通常为64KB。
• 广播机制：如果一个warp内的所有线程访问同一个常量地址，硬件会广播给所有线程，只产生一次内存事务。

常量内存最适合存储不会变化的参数。比如卷积核的权重、查找表、数学常数等。

使用常量内存很简单：

__constant__ float filter[256];

// 在主机端初始化
cudaMemcpyToSymbol(filter, host_filter, sizeof(float) * 256);

// 在kernel中使用
__global__ void applyFilter(float *output, float *input) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    output[idx] = input[idx] * filter[0];  // 所有线程读同一个地址
}

3.6 纹理内存（Texture Memory）

纹理内存最初是为图形渲染设计的。它有一些独特的特性：

• 硬件插值：支持线性插值、双线性插值等。做图像处理时特别有用。
• 边界处理：支持多种寻址模式（clamp、wrap、mirror等）。
• 空间局部性优化：纹理缓存针对2D空间局部性做了优化。如果你的数据访问模式是随机的，纹理内存可能比全局内存快。

我记得在做一个光线追踪项目时，场景中的纹理数据访问模式非常不规则。用全局内存读，缓存命中率很低。换成纹理内存后，性能提升了30%左右。

使用纹理内存的示例：

// 定义纹理对象
texture<float, 2, cudaReadModeElementType> tex;

// 绑定纹理
cudaArray *array;
cudaMallocArray(&array, &desc, width, height);
cudaMemcpyToArray(array, 0, 0, host_data, size, cudaMemcpyHostToDevice);
cudaBindTextureToArray(tex, array);

// 在kernel中使用
__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);  // 支持非整数坐标
}

3.7 内存层次总结

说了这么多，我整理了一个表格，方便你对比：

内存类型	位置	延迟	容量	作用域	缓存
寄存器	芯片内	~1 cycle	每个线程最多255个	单个线程	无
共享内存	芯片内	~5 cycles	每个SM 48-96KB	线程块	无
常量内存	显存	~400 cycles（缓存命中时~5 cycles）	64KB	全局	常量缓存
纹理内存	显存	~400 cycles（缓存命中时~5 cycles）	取决于显存	全局	纹理缓存
本地内存	显存	~400 cycles	每个线程最多512KB	单个线程	L1/L2
全局内存	显存	~400-800 cycles	最大（GB级别）	全局	L1/L2

最后说一句：没有银弹。每种内存都有它的适用场景。我见过有人为了用共享内存而用共享内存，结果性能反而更差。关键是要理解你的数据访问模式，然后选择最合适的存储层次。

下一章我们会深入讨论内存访问模式与性能优化。到时候我会分享一些实际项目中的优化案例，包括怎么用NVIDIA Nsight工具分析内存访问效率。敬请期待。