CUDA内存管理与优化：全局内存、共享内存、常量内存、纹理内存、寄存器

大家好，我是你们的老朋友。今天咱们聊聊CUDA里最核心、也最让人头疼的话题——内存管理。说实话，我见过太多人GPU算力利用率上不去，十有八九是内存访问出了问题。你想想看，GPU计算单元跑得飞快，但数据如果卡在路上，那再强的算力也是白搭。

这一章，我会带大家把CUDA的几种内存类型挨个捋一遍。每种内存有什么特点、什么时候用、怎么用才能榨干性能，我都会结合自己的实战经验来讲。嗯，咱们开始吧。

全局内存：最常用，也最容易踩坑

全局内存是GPU里容量最大的存储空间，所有线程都能访问。但它的延迟也最高——大概400-800个时钟周期。我在项目中遇到过不少新手，上来就把所有数据丢进全局内存，结果性能惨不忍睹。

内存合并访问是全局内存优化的第一要务。什么意思呢？就是让同一个warp里的32个线程，访问连续的内存地址。GPU硬件会把多个小请求合并成一次大的内存事务，效率能提升好几倍。

看个例子：

// 好的写法：合并访问
__global__ void good_kernel(float *data, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        float val = data[idx];  // 线程0访问data[0]，线程1访问data[1]...
        // 处理数据...
    }
}

// 坏的写法：非合并访问
__global__ void bad_kernel(float *data, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        float val = data[idx * 32];  // 线程0访问data[0]，线程1访问data[32]...
        // 性能会差很多
    }
}

我个人习惯在写kernel之前，先画一下线程和数据的映射关系。确保每个warp访问的地址是连续的，这能省下不少调试时间。

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

共享内存是GPU芯片上的SRAM，延迟只有几个时钟周期。但它有两个限制：容量小（通常48KB-96KB），而且只在同一个block内可见。

共享内存最大的价值在于数据复用。比如矩阵乘法，每个元素会被多个线程用到。如果每次都从全局内存读，那带宽根本不够用。把数据先搬到共享内存里，让block内的线程共享，性能能翻好几倍。

Bank Conflict：共享内存的隐形杀手

共享内存被分成32个bank（对应32个线程）。如果同一个warp里的多个线程访问同一个bank的不同地址，就会发生bank conflict，导致访问串行化。

举个例子：

// 有bank conflict的写法
__shared__ float smem[32][32];
int tid = threadIdx.x;
float val = smem[tid][tid];  // 所有线程访问同一列，都在同一个bank

// 无bank conflict的写法
float val = smem[tid][0];    // 每个线程访问不同行，分布在不同的bank

我曾经在一个图像处理项目里，因为没注意bank conflict，共享内存的带宽利用率只有理论值的1/4。后来改成按列访问，性能直接提升了3倍。嗯，这个坑我踩过，你们就别再踩了。

常量内存：只读数据的福音

常量内存有专门的缓存，延迟低，而且支持广播——如果一个warp里的所有线程访问同一个常量地址，硬件只会发一次请求。适合存放查找表、系数矩阵这类只读数据。

用法很简单：

__constant__ float coeff[256];

// 在host端初始化
cudaMemcpyToSymbol(coeff, host_coeff, sizeof(float) * 256);

// 在kernel里直接使用
__global__ void kernel(float *out) {
    int idx = threadIdx.x;
    out[idx] = coeff[idx % 256] * input[idx];
}

我个人建议，只要数据是只读的、且所有线程访问相同地址，优先考虑常量内存。但注意，常量内存只有64KB，别塞太多东西。

纹理内存：二维访问的利器

纹理内存专门为空间局部性优化过的。如果你处理的是图像、矩阵这类二维数据，纹理内存能利用硬件缓存，让相邻像素的访问更快。它还支持硬件插值，做图像缩放时特别方便。

不过说实话，纹理内存在通用计算里用得不多。除非你明确需要空间局部性优化，否则用全局内存+共享内存的组合就够了。

寄存器：最快的存储，但数量有限

寄存器是每个线程私有的，延迟为0。但每个线程能用的寄存器数量有限（通常255个）。如果kernel里变量太多，编译器会把多余的变量“溢出”到本地内存——那性能就崩了。

内存带宽优化策略总结

说了这么多，我给大家整理一个实用的检查清单：

优化目标	怎么做	我踩过的坑
全局内存合并访问	相邻线程访问相邻地址	曾经用列优先存储，结果带宽利用率不到10%
避免Bank Conflict	让不同线程访问不同bank	共享内存按列访问，忘了加padding
数据复用	用共享内存缓存重复数据	矩阵乘法没做分块，全局内存读爆了
常量内存广播	所有线程读同一个地址	把查找表放全局内存，浪费了广播特性
寄存器优化	控制变量数量，避免溢出	一个kernel用了300个寄存器，occupancy掉到25%

知识体系总览

下面这张图，是我自己画的内存体系结构。你看一眼就能明白各种内存的位置和关系：

你看，从寄存器到全局内存，延迟差了上百倍。所以优化的核心思路就一句话：把频繁访问的数据尽量往“上面”放。

好了，这一章的内容就到这里。内存优化是个细活，需要结合profiler反复调。下一章咱们会讲NCCL的通信优化，到时候再聊。