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

聊到CUDA编程，内存层次结构是绕不开的核心话题。我见过太多新手，一上来就写核函数，结果性能惨不忍睹。说白了，问题多半出在内存访问上。

GPU的内存不是铁板一块。它分了好几个层级，每个层级的速度、容量、用途都不一样。你想想看，如果能把数据放在最合适的地方，性能翻个几倍不是梦。

寄存器：最快但最稀缺

寄存器是离计算单元最近的地方。访问它只需要一个时钟周期，几乎没有延迟。我个人习惯把最频繁使用的变量放在寄存器里。

我在项目中遇到过这样一个坑：核函数里定义了太多局部变量，结果寄存器不够用。编译器默默地把一部分变量挪到了本地内存，性能掉了30%。后来我强制限制了每个线程的寄存器数量（用__launch_bounds__），问题才解决。

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

本地内存：寄存器的“备胎”

本地内存听起来像“本地”的，其实它住在全局内存里。访问延迟高得吓人。为什么会这样？因为寄存器不够用了，编译器只能把数据塞到全局内存里，用本地内存的“马甲”包装一下。

我曾经调试过一个核函数，发现性能瓶颈就在本地内存访问上。检查代码后发现，有个大数组定义在核函数内部，编译器没法把它全塞进寄存器。解决方案很简单：要么拆分成多个小变量，要么用共享内存。

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

共享内存是块宝。它在芯片上，离SM很近，延迟只有几十个时钟周期。同一个block里的线程都能访问它。说白了，这就是线程间协作的“会议室”。

共享内存的典型用途有两个：

• 数据复用：多个线程需要读取同一份数据时，先加载到共享内存，避免重复访问全局内存
• 线程间通信：block内的线程通过共享内存交换数据，比用全局内存快得多

// 共享内存示例：矩阵分块乘法
__global__ void matMulShared(const float* A, const float* B, float* C, int N) {
    __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
    
    int bx = blockIdx.x, by = blockIdx.y;
    int tx = threadIdx.x, ty = threadIdx.y;
    
    float sum = 0.0f;
    for (int k = 0; k < N / BLOCK_SIZE; ++k) {
        // 协作加载数据到共享内存
        As[ty][tx] = A[by * BLOCK_SIZE + ty][k * BLOCK_SIZE + tx];
        Bs[ty][tx] = B[k * BLOCK_SIZE + ty][bx * BLOCK_SIZE + tx];
        __syncthreads();
        
        for (int i = 0; i < BLOCK_SIZE; ++i)
            sum += As[ty][i] * Bs[i][tx];
        __syncthreads();
    }
    C[by * BLOCK_SIZE + ty][bx * BLOCK_SIZE + tx] = sum;
}

全局内存：容量大但速度慢

全局内存是GPU的主存，容量最大（几GB到几十GB），但延迟也最高（几百个时钟周期）。所有线程都能读写它，但访问效率取决于合并访问。

什么叫合并访问？就是相邻的线程访问相邻的内存地址。GPU会把多个线程的访问请求合并成一次大的内存事务。你想想看，如果每个线程都乱跳地址，那效率得多低？

访问模式	带宽利用率	说明
顺序合并访问	接近100%	线程0访问地址0，线程1访问地址1，以此类推
随机访问	10%-30%	每个线程访问的地址毫无规律
广播访问	100%	所有线程访问同一个地址

常量内存：只读且带缓存

常量内存是全局内存的一个特殊区域，有专门的缓存（常量缓存）。它适合存储那些所有线程都读取的、不变的参数。比如物理常数、查找表等。

常量内存的访问延迟很低，前提是所有线程访问同一个地址。如果每个线程访问不同的常量地址，性能反而比全局内存还差。

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

// 主机端赋值
cudaMemcpyToSymbol(params, hostParams, sizeof(float) * 256);

// 核函数中使用
__global__ void kernel() {
    float p = params[threadIdx.x % 256];  // 所有线程访问同一个地址？不一定！
}

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

纹理内存最初是为图形渲染设计的，但它对某些计算场景特别有用。它也有缓存，而且针对二维和三维的空间局部性做了优化。

什么时候用纹理内存？

• 数据访问有空间局部性（相邻线程访问相邻的纹理坐标）
• 需要硬件插值功能（比如图像处理中的双线性插值）
• 数据是只读的，且访问模式不规则

我记得有个项目是做图像滤波，用纹理内存比用全局内存快了3倍。因为纹理缓存在处理二维空间访问时，命中率特别高。

// 纹理内存声明
texture<float, 2, cudaReadModeElementType> texInput;

// 核函数中使用
__global__ void filterKernel(float* output, int width, int height) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    
    // 使用tex2D读取纹理内存
    float val = tex2D(texInput, x, y);
    output[y * width + x] = val * 2.0f;
}

一张图看懂内存层次

下面这张SVG图，把整个内存层次结构串起来了。你可以看到每个层级的位置、速度、容量和可见范围。

嗯，这张图把整个内存层次串起来了。从顶部的寄存器到底部的全局内存，速度越来越慢，容量越来越大。你在写核函数时，心里要时刻装着这张图。

好了，内存层次结构就聊到这儿。记住一句话：没有最好的内存，只有最合适的内存。每个层级都有自己的用武之地，关键看你如何搭配使用。