3. CUDA内存模型：全局内存、共享内存、寄存器、本地内存、常量内存、纹理内存

说到CUDA编程，内存模型这块儿，我敢说它是决定你程序性能的「命门」。很多人刚开始学CUDA，把精力全放在并行算法设计上，结果一跑起来，发现比CPU还慢。嗯，十有八九是内存访问出了问题。

我个人习惯，在写任何CUDA kernel之前，先画一张内存访问的草图。你想想看，数据从哪儿来，到哪儿去，中间经过哪些缓存，每一步的延迟是多少——心里有数了，代码才不会跑偏。

3.1 全局内存——最常用也最慢

全局内存，说白了就是GPU上的「主存」。CPU端用 cudaMalloc 分配的就是它。所有线程都能访问，容量最大，但延迟也最高——大概400到800个时钟周期。

我在项目中遇到过一个问题：一个矩阵乘法kernel，理论计算量不大，但跑起来就是慢。后来用 nvprof 一分析，发现全局内存访问的吞吐量只有理论峰值的10%。为什么？因为访问模式不对。

// ❌ 非合并访问——性能杀手
__global__ void bad_access(float *A, int N) {
    int tid = threadIdx.x;
    for (int i = 0; i < N; i++) {
        A[i * blockDim.x + tid] = 1.0f;  // 跨步访问
    }
}

// ✅ 合并访问——高效
__global__ void good_access(float *A, int N) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid < N) {
        A[tid] = 1.0f;  // 连续访问
    }
}

3.2 共享内存——线程块的「高速缓存」

共享内存是片上的，延迟只有5-10个周期，比全局内存快两个数量级。但它有两个限制：一是每个线程块只有48KB（有些架构可以配置到96KB），二是数据只在同一个线程块内可见。

我经常把共享内存比喻成「团队协作的白板」——同一个block里的线程可以往上面写东西，也可以读别人写的东西。但不同block之间，不好意思，各用各的白板。

// 共享内存做矩阵分块——经典用法
__global__ void matmul_shared(float *A, float *B, float *C, int N) {
    __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
    
    int row = blockIdx.y * BLOCK_SIZE + threadIdx.y;
    int col = blockIdx.x * BLOCK_SIZE + threadIdx.x;
    float sum = 0.0f;
    
    for (int tile = 0; tile < N / BLOCK_SIZE; tile++) {
        // 协作加载到共享内存
        As[threadIdx.y][threadIdx.x] = A[row * N + tile * BLOCK_SIZE + threadIdx.x];
        Bs[threadIdx.y][threadIdx.x] = B[(tile * BLOCK_SIZE + threadIdx.y) * N + col];
        __syncthreads();  // 同步！确保所有线程加载完成
        
        for (int k = 0; k < BLOCK_SIZE; k++) {
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }
        __syncthreads();  // 同步！防止覆盖数据
    }
    C[row * N + col] = sum;
}

3.3 寄存器——最快的存储，但很稀缺

寄存器是每个线程私有的，访问延迟为0——说白了就是一个时钟周期内就能拿到数据。但每个SM的寄存器总数是有限的，比如在A100上是65536个。如果每个线程用太多寄存器，就会限制SM上同时运行的线程块数量，这叫「寄存器压力」。

我一般会关注编译器的输出信息：

// 编译时加上 --ptxas-options=-v 可以看到寄存器使用量
// nvcc -arch=sm_80 --ptxas-options=-v mykernel.cu

// 如果发现寄存器使用量过高（比如超过64），可以考虑：
// 1. 用 __launch_bounds__ 提示编译器
// 2. 把一些局部变量改成 volatile 或拆分成更小的kernel
__global__ __launch_bounds__(256, 4) void my_kernel() {
    // 限制每个block最多256个线程，每个SM至少4个block
}

3.4 本地内存——寄存器的「备胎」

本地内存听起来像是「本地」的，其实它物理上就在全局内存里。编译器在两种情况下会把变量放到本地内存：一是寄存器不够用了（溢出），二是有些数组索引在编译期无法确定。

说白了，本地内存就是个「兜底」方案。你写代码时感觉不到它的存在，但性能分析时如果发现 local_load/store 指令很多，就要警惕了。

3.5 常量内存——只读且带广播

常量内存只有64KB，但它的厉害之处在于：当一个warp里的所有线程访问同一个常量地址时，硬件只需要一次读取，然后广播给所有线程。延迟只有10-20个周期（如果缓存命中）。

// 常量内存适合存放查找表、系数等只读数据
__constant__ float filter[256];

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

// kernel中直接使用——所有线程共享同一份缓存
__global__ void apply_filter(float *in, float *out, int N) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        out[idx] = in[idx] * filter[threadIdx.x & 0xFF];
    }
}

3.6 纹理内存——为空间局部性而生

纹理内存最初是为图形学设计的，但它的一些特性在通用计算中也很有用：支持硬件插值、边界处理、以及针对2D空间局部性的缓存优化。

我在做图像处理时用过纹理内存。比如做图像卷积，相邻像素的访问模式天然就有空间局部性，纹理缓存的命中率很高。不过说实话，现在统一内存架构越来越成熟，纹理内存的使用场景在逐渐减少。

// 纹理内存的典型用法
texture<float, 2, cudaReadModeElementType> tex;

// 绑定纹理
cudaArray *cuArray;
cudaMallocArray(&cuArray, &channelDesc, width, height);
cudaMemcpyToArray(cuArray, 0, 0, hostData, size, cudaMemcpyHostToDevice);
cudaBindTextureToArray(tex, cuArray);

// kernel中通过tex2D访问
__global__ void kernel(float *out, int width, int height) {
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    if (x < width && y < height) {
        out[y * width + x] = tex2D(tex, x + 0.5f, y + 0.5f);
    }
}

3.7 如何选择——我的决策框架

每次写kernel，我都会问自己几个问题：

1. 数据是只读还是读写？ 只读且warp内访问一致 → 常量内存；只读且空间局部性强 → 纹理内存；读写 → 全局内存或共享内存。
2. 数据是否被同一个block的多个线程复用？ 是 → 用共享内存做缓存；否 → 直接全局内存访问，注意合并。
3. 每个线程的私有数据量有多大？ 小（几个变量）→ 寄存器；大（数组或结构体）→ 小心寄存器溢出，考虑拆分kernel。

好了，CUDA的内存模型就聊到这儿。每种内存都有它的脾气，摸透了，你的kernel性能就不会差。记住一句话：让数据尽量靠近计算，让访问模式尽量规整。