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

说到CUDA编程，内存管理绝对是绕不开的核心话题。我刚开始接触CUDA时，总觉得把数据扔到GPU上就能自动跑得快，结果被现实狠狠教育了一顿。后来才明白，不理解内存层次结构，就像开着跑车却走错了路——再好的硬件也白搭。

GPU的内存不是一块铁板，而是一个金字塔结构。越靠近计算单元的内存越快，但容量也越小。咱们今天就把这六种内存掰开揉碎了讲清楚。

核心要点：内存访问模式决定了你的核函数能跑多快。选对内存类型，性能翻倍不是梦。

1. 全局内存（Global Memory）

全局内存是GPU上最大的存储空间，所有线程都能访问。说白了，它就是GPU的“主存”。容量大，但延迟也高——大概400-800个时钟周期。

我个人习惯把全局内存比作“仓库”。你需要什么数据，得亲自去仓库搬。搬一次很慢，但仓库够大，啥都能存。

// 全局内存声明示例
__global__ void kernel(float *d_data) {
    // d_data 指向的就是全局内存
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    d_data[idx] = d_data[idx] * 2.0f;
}

注意：全局内存的访问要尽量合并（coalesced）。相邻线程访问相邻地址，才能发挥带宽优势。我曾经见过一个项目，就因为没注意合并访问，性能直接掉了5倍。

2. 共享内存（Shared Memory）

共享内存是块上的“小食堂”。同一个块里的线程都能访问，速度比全局内存快得多——大概1-2个时钟周期。但容量有限，每个SM通常只有几十KB到上百KB。

我在项目中遇到过这样一个场景：矩阵乘法需要反复读取子矩阵。如果每次都去全局内存拿，那效率简直没法看。用共享内存做缓存，速度直接起飞。

// 共享内存使用示例
__global__ void shared_mem_kernel(float *d_in, float *d_out) {
    __shared__ float s_data[256];  // 声明共享内存
    
    int tid = threadIdx.x;
    s_data[tid] = d_in[tid];  // 从全局内存加载到共享内存
    __syncthreads();           // 同步，确保所有线程加载完成
    
    // 在共享内存上做计算
    s_data[tid] = s_data[tid] * 2.0f;
    __syncthreads();
    
    d_out[tid] = s_data[tid];  // 写回全局内存
}

避坑指南：我曾经因为忘记加 __syncthreads()，导致线程间数据不一致，调试了整整两天。记住：共享内存的读写必须同步，否则会出现竞态条件。

3. 寄存器（Registers）

寄存器是GPU上最快的内存，没有之一。每个线程独享，延迟几乎为零。但数量极其有限——每个线程通常只有几十到几百个寄存器。

你想想看，如果每个线程用太多寄存器，SM能同时调度的线程数就会减少，这叫“寄存器压力”。我见过有人为了省事，把所有变量都声明成寄存器，结果占用率掉到30%，性能反而更差。

// 寄存器自动管理，编译器会优化
__global__ void reg_kernel() {
    float a = 1.0f;  // 编译器会尽量放在寄存器
    float b = 2.0f;
    float c = a + b; // 寄存器操作，极快
}

注意：别手动声明 register 关键字，编译器比你聪明。它会自动决定哪些变量放寄存器，哪些放本地内存。

4. 本地内存（Local Memory）

本地内存听起来像“本地”的，其实它和全局内存一样慢。为什么叫本地？因为它对每个线程是私有的，其他线程看不到。

什么时候会用本地内存？当寄存器不够用的时候。比如你声明了一个大数组，或者编译器觉得某个变量不适合放寄存器，就会自动“溢出”到本地内存。

// 这种情况容易触发本地内存
__global__ void local_mem_kernel() {
    float big_array[100];  // 寄存器放不下，会放到本地内存
    // 访问 big_array 实际上是在访问全局内存
}

核心要点：本地内存是“不得已的选择”。如果发现核函数性能异常，可以用 --ptxas-options=-v 编译选项查看本地内存使用量。

5. 常量内存（Constant Memory）

常量内存是GPU上一块特殊的内存区域，容量只有64KB。它被所有线程共享，但只能读不能写。为什么特殊？因为它有专门的缓存机制，当所有线程访问同一个地址时，速度极快。

我记得有个项目需要做查找表（LUT），把所有线程都用同一个系数表。用常量内存后，性能提升了30%。

// 常量内存声明
__constant__ float c_coeff[256];  // 在主机端用 cudaMemcpyToSymbol 初始化

__global__ void const_kernel(float *d_out) {
    int idx = threadIdx.x;
    d_out[idx] = d_in[idx] * c_coeff[idx % 256];  // 读取常量内存
}

避坑指南：常量内存最适合“所有线程读同一个地址”的场景。如果每个线程读不同地址，反而比全局内存还慢。我曾经踩过这个坑，切记。

6. 纹理内存（Texture Memory）

纹理内存最初是为图形处理设计的，但在通用计算中也有妙用。它支持硬件插值、边界处理，还有专门的缓存。对于二维空间局部性好的数据，纹理内存表现很出色。

说实话，纹理内存现在用得少了。因为统一内存和更智能的缓存机制已经能覆盖大部分场景。但在某些特定场景——比如图像处理、随机访问模式——它依然有优势。

// 纹理内存使用示例（简化版）
texture<float, 2, cudaReadModeElementType> tex;

__global__ void texture_kernel(float *d_out) {
    int x = threadIdx.x;
    int y = threadIdx.y;
    // 通过纹理读取，自动处理边界
    d_out[y * blockDim.x + x] = tex2D(tex, x + 0.5f, y + 0.5f);
}

内存层次结构总览

为了让你看得更清楚，我画了一张图。这张图展示了GPU内存的金字塔结构：

如何选择合适的内存？

这个问题没有标准答案，但我可以给你一些经验法则：

频繁访问的标量变量 → 交给寄存器，编译器会自动处理
块内线程需要共享数据 → 用共享内存，记得加同步
所有线程读同一个只读数据 → 常量内存，64KB够用
二维空间局部性好的数据 → 纹理内存，或者用共享内存手动优化
大容量、随机访问 → 全局内存，注意合并访问
寄存器溢出 → 尽量避免，如果发生就重构代码

总结一下：内存层次结构是CUDA性能优化的基石。我见过太多人一上来就写核函数，结果被内存访问拖垮。花点时间理解这六种内存，绝对值得。

嗯，今天就聊到这儿。记住：选对内存类型，你的核函数才能跑出真正的性能。

📚 课程目录

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

1. 全局内存（Global Memory）

2. 共享内存（Shared Memory）

3. 寄存器（Registers）

4. 本地内存（Local Memory）

5. 常量内存（Constant Memory）

6. 纹理内存（Texture Memory）

内存层次结构总览

如何选择合适的内存？