3、CUDA内存模型：全局内存、共享内存、寄存器、本地内存、常量内存、纹理内存的特点与使用场景

说到CUDA编程，内存模型这块儿是绕不开的坎儿。我刚开始学的时候，总觉得不就是几个存储空间嘛，有啥好纠结的？结果一上手写kernel，性能惨不忍睹。后来才明白，不理解内存模型，你写的CUDA代码就是碰运气。

今天咱们就把这六种内存掰开揉碎了讲清楚。每种内存什么脾气、什么时候用、怎么用，我都会结合自己的实战经验来说。

3.1 全局内存（Global Memory）

全局内存是CUDA里最大的存储空间，也是所有线程都能访问的。说白了，它就是GPU上的DRAM，容量大（几个GB到几十个GB），但延迟也高——一次访问大概要400到800个时钟周期。

我有个习惯：能用共享内存解决的问题，绝不用全局内存。但有些数据是躲不开的，比如输入数据、输出结果，这些必须放在全局内存里。

// ❌ 非合并访问：性能差
__global__ void bad_access(float *A, float *B) {
    int tid = threadIdx.x;
    B[tid] = A[tid * 32];  // 每个线程间隔32个元素
}

// ✅ 合并访问：性能好
__global__ void good_access(float *A, float *B) {
    int tid = threadIdx.x;
    B[tid] = A[tid];  // 连续访问
}

3.2 共享内存（Shared Memory）

共享内存是我最喜欢的存储类型。它位于芯片上，延迟只有几个周期，比全局内存快了两个数量级。但容量很小，每个SM上通常只有几十KB到一百多KB。

共享内存的作用是什么？让同一个block里的线程协作。比如做矩阵乘法时，先把数据从全局内存搬到共享内存，然后所有线程从共享内存里读数据，这样就能避免反复访问慢速的全局内存。

// 共享内存声明：动态分配
__global__ void kernel() {
    extern __shared__ float shared_data[];
    // 使用 shared_data
}

// 静态分配
__global__ void kernel2() {
    __shared__ float cache[256];
}

3.3 寄存器（Register）

寄存器是速度最快的存储，没有之一。每个线程都有自己的私有寄存器，编译器会自动分配。但寄存器数量有限——每个SM的寄存器总数是固定的（比如65536个），如果每个线程用太多，就会降低occupancy（驻留线程数）。

我个人的经验是：别让编译器替你操心寄存器分配。有时候你写了个很复杂的kernel，编译器一优化，寄存器用量飙升，结果occupancy掉到谷底。这时候可以用 __launch_bounds__ 来限制每个线程的寄存器数量。

// 限制每个线程最多使用32个寄存器
__global__ __launch_bounds__(256, 4) void my_kernel() {
    // 你的代码
}

3.4 本地内存（Local Memory）

本地内存听起来像是"本地"的，其实它就是个幌子。它物理上位于全局内存里，只是每个线程私有的。什么时候会用？寄存器不够用的时候，编译器会自动把变量放到本地内存里。

嗯，这里要注意：本地内存的访问延迟和全局内存一样高。所以如果你发现kernel性能不对劲，可以检查一下PTX代码，看看有没有大量的 st.local 或 ld.local 指令。如果有，说明寄存器溢出了。

3.5 常量内存（Constant Memory）

常量内存也是全局内存的一部分，但它有专门的缓存（constant cache），而且有硬件广播机制。什么意思呢？如果一个warp里的所有线程都访问同一个常量地址，那只需要一次内存访问，然后广播给所有线程。

常量内存适合存什么？只读的、所有线程都访问相同地址的数据。比如神经网络的权重、查找表、滤波器系数等。

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

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

// kernel中使用
__global__ void inference(float *input, float *output) {
    int idx = threadIdx.x;
    output[idx] = input[idx] * weights[idx];  // 所有线程访问相同的weights
}

3.6 纹理内存（Texture Memory）

纹理内存最初是为图形渲染设计的，但用在通用计算里也有奇效。它有几个特点：

• 有硬件缓存：对空间局部性好的访问模式很友好
• 支持插值：硬件自动做双线性/三线性插值
• 支持边界处理：可以设置边界外的返回值
• 对2D访问优化：访问二维数据时，比全局内存快很多

我一般在图像处理、计算摄影、或者需要随机访问二维数据的场景下用纹理内存。比如做图像卷积时，用纹理内存比全局内存快30%以上。

// 声明纹理引用
texture<float, 2, cudaReadModeElementType> tex;

// 绑定纹理
cudaBindTexture2D(NULL, tex, dev_data, &tex_channel_desc, width, height, pitch);

// kernel中使用
__global__ void filter_kernel(float *output) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    float val = tex2D(tex, x + 0.5f, y + 0.5f);  // 硬件插值
    output[y * width + x] = val;
}

3.7 总结对比

内存类型	位置	延迟	容量	作用域	典型用途
寄存器	片上	~1周期	每个线程最多255个	线程私有	局部变量、中间结果
共享内存	片上	~5周期	每个SM 48~164KB	Block内共享	线程协作、数据重用
本地内存	片外(DRAM)	~400周期	同全局内存	线程私有	寄存器溢出
全局内存	片外(DRAM)	~400周期	几GB~几十GB	所有线程	输入输出数据
常量内存	片外(DRAM)+缓存	~100周期(缓存命中)	64KB	所有线程(只读)	权重、查找表
纹理内存	片外(DRAM)+缓存	~100周期(缓存命中)	同全局内存	所有线程(只读)	图像、2D数据

好了，CUDA内存模型就讲到这里。每种内存都有自己的脾气，摸透了它们，你写的kernel才能又快又稳。记住一句话：让数据离计算越近越好。