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

说到CUDA编程，内存管理绝对是绕不开的核心话题。我刚开始接触GPU计算时，总觉得把数据扔到显存里就完事了。结果呢？性能惨不忍睹。后来我才明白——不理解内存层次结构，你写的核函数就是在浪费GPU的生命。

说白了，GPU的内存系统就像一座金字塔。越靠近计算单元，速度越快，但容量越小。咱们今天就把这六种内存掰开揉碎了讲清楚。

3.1 全局内存：最基础也最慢

全局内存是GPU中容量最大的存储区域，通常有数GB。所有线程都能读写它，但代价是极高的延迟——大约400到800个时钟周期。

我在项目中遇到过这样的情况：一个简单的向量加法，直接读写全局内存，带宽利用率只有20%左右。为什么？因为全局内存访问需要满足合并访问的条件。

核心原则：相邻线程访问相邻地址，才能实现带宽最大化。

举个例子，假设每个线程读取一个float（4字节），如果线程0读地址0，线程1读地址4，线程2读地址8……这样32个线程的访问可以合并成一次128字节的事务。反之，如果线程0读地址0，线程1读地址100，那就得拆成多次事务，带宽直接打骨折。

// 好的合并访问：线程i访问data[i]
__global__ void good_access(float *data, int N) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        data[idx] = data[idx] * 2.0f;  // 合并访问
    }
}

// 差的访问模式：线程i访问data[gridDim.x * idx]
__global__ void bad_access(float *data, int N) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        data[gridDim.x * idx] = data[gridDim.x * idx] * 2.0f;  // 跨步访问
    }
}

注意：全局内存的读写不会自动缓存（L1/L2缓存存在，但行为不可控）。如果你需要频繁重用数据，别指望全局内存的缓存——用共享内存才是正道。

3.2 共享内存：可编程的缓存

共享内存是块内线程共享的存储空间，延迟大约30个周期，比全局内存快一个数量级。每个SM（流式多处理器）通常有48KB到96KB的共享内存，具体取决于架构。

我个人习惯把共享内存当作手动管理的缓存。你想想看，如果每个线程都要从全局内存读同一份数据，为什么不先把它搬到共享内存里，然后大家直接从共享内存读呢？

// 矩阵分块乘法：使用共享内存减少全局内存访问
__global__ void matmul_shared(float *A, float *B, float *C, int N) {
    __shared__ float As[TILE_SIZE][TILE_SIZE];
    __shared__ float Bs[TILE_SIZE][TILE_SIZE];
    
    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    float sum = 0.0f;
    
    for (int t = 0; t < N / TILE_SIZE; t++) {
        // 协作加载数据到共享内存
        As[threadIdx.y][threadIdx.x] = A[row * N + t * TILE_SIZE + threadIdx.x];
        Bs[threadIdx.y][threadIdx.x] = B[(t * TILE_SIZE + threadIdx.y) * N + col];
        __syncthreads();  // 确保所有线程加载完成
        
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }
        __syncthreads();  // 防止下一轮迭代覆盖数据
    }
    
    C[row * N + col] = sum;
}

避坑指南：我曾经在项目里忘了加__syncthreads()，结果数据竞争导致计算结果时对时错。调试了整整两天才发现问题。记住：共享内存的读写必须用同步原语来保证顺序。

3.3 寄存器：最快的存储

寄存器是每个线程私有的，访问延迟为零。每个线程最多能用的寄存器数量有限——通常是255个（取决于架构和编译器设置）。

如果每个线程使用的寄存器超过限制，会发生什么？寄存器溢出。溢出的数据会被放到本地内存里，而本地内存实际上是在全局内存中分配的。这一下，性能就崩了。

我建议你在编译时加上--ptxas-options=-v参数，看看每个核函数用了多少寄存器。如果超过32个，就得考虑优化了。

// 查看寄存器使用量
// nvcc --ptxas-options=-v mykernel.cu
// 输出示例：
//   注册使用量: 24 registers, 0 bytes stack, 0 bytes spill stores, 0 bytes spill loads

经验之谈：寄存器不是越多越好。占用率（occupancy）和寄存器数量是矛盾的。每个SM的寄存器总数固定（比如65536个），如果每个线程用太多寄存器，能同时驻留的线程数就少了。我一般控制在32个寄存器以内，保证至少50%的占用率。

3.4 本地内存：寄存器的替补

本地内存听起来像是"本地"的，但实际上它存储在全局内存中，延迟很高。编译器会自动把放不下的局部变量、大型结构体、数组等 spill 到本地内存。

你想想看，如果一个线程声明了一个很大的局部数组：

__global__ void kernel() {
    float big_array[1024];  // 这很可能被放到本地内存
    // ...
}

这个数组大概率会被分配到本地内存。每次访问都要走全局内存的延迟，性能可想而知。

注意：本地内存的访问模式也要注意合并。虽然每个线程的本地内存是私有的，但硬件会尽量合并同一warp内线程的本地内存访问。不过，这远不如直接使用共享内存来得高效。

3.5 常量内存：只读且缓存友好

常量内存只有64KB，但它的缓存机制非常高效。当warp内所有线程访问同一个常量地址时，只需要一次内存事务。如果访问不同地址，那就串行化了——性能会下降。

我在项目中常用常量内存来存储查找表、滤波器系数等只读数据。比如图像处理中的高斯核系数，放在常量内存里，所有线程共享，缓存命中率极高。

// 声明常量内存
__constant__ float filter[256];

// 主机端复制数据
cudaMemcpyToSymbol(filter, host_filter, sizeof(float) * 256);

// 核函数中使用
__global__ void apply_filter(float *input, float *output, int N) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        float val = input[idx];
        // 所有线程访问同一个filter地址，缓存友好
        output[idx] = val * filter[0] + val * filter[1];
    }
}

小技巧：常量内存的读取延迟大约和共享内存相当（如果缓存命中）。但要注意，常量内存的更新需要同步——所有正在执行的核函数必须完成，才能更新常量内存。所以它适合"一次设置，多次使用"的场景。

3.6 纹理内存：空间局部性的王者

纹理内存最初是为图形学设计的，但它对某些计算模式有奇效。纹理内存有专门的缓存，针对空间局部性做了优化——如果相邻线程访问相邻的纹理坐标，缓存命中率极高。

另外，纹理内存还支持硬件插值、边界处理等功能。对于图像处理、科学计算中的网格数据访问，纹理内存往往比全局内存快得多。

// 使用纹理内存读取2D数据
texture<float, 2, cudaReadModeElementType> tex;

__global__ void texture_kernel(float *output, 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) {
        // 使用tex2D读取纹理，自动处理边界
        output[y * width + x] = tex2D(tex, x + 0.5f, y + 0.5f);
    }
}

什么时候用纹理内存？我个人总结三条：

数据访问模式有空间局部性（相邻线程访问相邻地址）
需要硬件插值或边界处理
只读数据且访问模式不规则（纹理缓存的策略比L1缓存更智能）

3.7 如何选择？一张表说清楚

内存类型	作用域	生命周期	延迟	容量	典型用途
寄存器	线程私有	核函数执行期间	0 cycles	有限（每个线程最多255个）	局部变量、中间结果
本地内存	线程私有	核函数执行期间	高（同全局内存）	取决于寄存器溢出	寄存器溢出的数据
共享内存	块内共享	块执行期间	约30 cycles	48KB/block（可配置）	块内数据重用、协作计算
全局内存	所有线程+主机	应用程序全程	400-800 cycles	GB级	主要数据存储
常量内存	所有线程只读	应用程序全程	约30 cycles（缓存命中）	64KB	查找表、系数
纹理内存	所有线程只读	应用程序全程	约100 cycles（缓存命中）	取决于设备	图像、网格数据

嗯，到这里你应该对CUDA的内存层次有了清晰的认识。记住一句话：让数据尽可能靠近计算单元。全局内存是最后的选择，共享内存是性能优化的利器，寄存器是免费的午餐但别贪多。至于常量内存和纹理内存，它们各有专长，用对了就是神兵利器。

我见过太多开发者一上来就写全局内存访问，然后抱怨GPU不够快。其实很多时候，换个内存类型，性能就能翻倍。你想想看，是不是这个理？

📚 课程目录

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

3.1 全局内存：最基础也最慢

3.2 共享内存：可编程的缓存

3.3 寄存器：最快的存储

3.4 本地内存：寄存器的替补

3.5 常量内存：只读且缓存友好

3.6 纹理内存：空间局部性的王者

3.7 如何选择？一张表说清楚