第三章：内存模型——全局内存、共享内存、寄存器、本地内存、常量内存、纹理内存

聊到CUDA编程，绕不开的就是内存模型。说实话，很多新手甚至老手，性能瓶颈都卡在内存访问上。我个人习惯把GPU的内存体系比作一个“分层仓库”，每一层速度不同、容量不同、用法也不同。搞懂了这些，你的算子优化才算真正入门。

3.1 全局内存（Global Memory）

全局内存是GPU里容量最大、但延迟也最高的存储空间。所有线程都能读写它，CPU也能通过PCIe把数据搬进来。说白了，它就是显存本身。

特点：

容量大（几GB到几十GB）
延迟高（几百个时钟周期）
所有线程可见
支持动态分配（cudaMalloc）

避坑指南：我曾经在项目里犯过一个低级错误——频繁在kernel里用全局内存做原子操作。结果性能惨不忍睹。后来改成先用共享内存做局部归约，再写回全局内存，速度直接翻了5倍。

访问全局内存时，有个关键概念叫“合并访问”（coalesced access）。什么意思呢？就是相邻线程最好访问相邻地址。你想想看，如果线程0读地址0，线程1读地址1，线程2读地址2……这样硬件可以把多次访问合并成一次，效率极高。反之，如果线程0读地址0，线程1读地址100，线程2读地址200……那就得三次独立访问，性能直接崩了。

// 合并访问示例（推荐）
__global__ void coalesced_kernel(float *data, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        data[idx] = data[idx] * 2.0f;  // 相邻线程访问相邻地址
    }
}

// 非合并访问示例（不推荐）
__global__ void non_coalesced_kernel(float *data, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        data[threadIdx.x * N + blockIdx.x] = data[threadIdx.x * N + blockIdx.x] * 2.0f;  // 跨步访问
    }
}

3.2 共享内存（Shared Memory）

共享内存是每个block私有的高速缓存，容量很小（通常几十KB），但延迟极低（几个时钟周期）。它最大的用处是让同一个block内的线程协作。

特点：

容量小（通常16KB-48KB可配置）
延迟低（约5-10个时钟周期）
同一block内线程共享
需要手动管理数据搬入搬出

我的经验：共享内存最经典的用法是做矩阵分块。比如矩阵乘法，把子块先加载到共享内存里，然后反复使用，避免反复访问全局内存。我做过一个测试，用共享内存优化后的矩阵乘法，比直接读全局内存快了将近10倍。

嗯，这里要注意一个坑：共享内存有bank conflict问题。共享内存被分成32个bank，如果多个线程同时访问同一个bank的不同地址，就会发生冲突，导致串行化。解决办法？让线程访问的地址错开，或者用padding技巧。

// 共享内存声明
__shared__ float shared_data[256];

// 典型用法：把全局内存数据加载到共享内存
__global__ void shared_memory_kernel(float *input, float *output, int N) {
    __shared__ float cache[256];
    int tid = threadIdx.x;
    int idx = blockIdx.x * blockDim.x + tid;
    
    if (idx < N) {
        cache[tid] = input[idx];  // 加载到共享内存
    }
    __syncthreads();  // 同步，确保所有线程都加载完毕
    
    // 在共享内存上做计算
    // ...
    
    if (idx < N) {
        output[idx] = cache[tid];  // 写回全局内存
    }
}

3.3 寄存器（Register）

寄存器是GPU里最快、最私有的存储空间。每个线程独享，延迟几乎为0。但数量极其有限——每个线程通常只有几十到几百个寄存器。

特点：

速度最快（1个时钟周期）
每个线程私有
数量有限（通常每个线程最多255个）
编译器自动管理

注意：寄存器用多了会 spill 到本地内存，性能会急剧下降。我曾经调试过一个kernel，寄存器用了300多个，结果编译器偷偷把一部分变量塞到了本地内存里，性能直接腰斩。后来我手动优化了代码，减少了中间变量，才把寄存器压到合理范围。

怎么控制寄存器使用量？可以用编译选项 --maxrregcount=32 限制每个线程最多用32个寄存器。但别设得太死，否则编译器会疯狂spill，反而更慢。

3.4 本地内存（Local Memory）

本地内存听起来像“本地”的，其实它物理上就在全局内存里。只是每个线程私有的全局内存而已。当寄存器不够用，或者访问数组索引不是编译期常量时，编译器会把变量放到本地内存里。

特点：

物理上位于全局内存
每个线程私有
延迟高（和全局内存一样）
编译器自动决定使用

为什么会用到本地内存？举个例子：你声明了一个大数组 int arr[100]，寄存器放不下，编译器就会把它放到本地内存。或者你用了 arr[threadIdx.x] 这种动态索引，也可能触发本地内存。

避坑指南：我曾经写过一个kernel，里面有个大结构体，结果编译器把它全塞进了本地内存。性能惨不忍睹。后来我把结构体拆开，用多个小数组代替，寄存器使用量降下来了，性能也恢复了。

3.5 常量内存（Constant Memory）

常量内存是只读的，容量很小（64KB），但有个神奇的特性——如果所有线程访问同一个地址，它比全局内存快得多。因为它有专门的缓存机制。

特点：

只读
容量小（64KB）
有缓存
适合存储常量参数

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

// 主机端赋值
cudaMemcpyToSymbol(const_params, host_params, sizeof(float) * 256);

// 设备端使用
__global__ void constant_kernel(float *output, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        output[idx] = output[idx] * const_params[0];  // 所有线程读同一个地址
    }
}

我个人习惯把卷积核的权重、归一化参数这类所有线程都要用的常量放到常量内存里。注意，如果每个线程访问不同的地址，常量内存反而比全局内存慢，因为缓存会失效。

3.6 纹理内存（Texture Memory）

纹理内存最初是为图形渲染设计的，但用在通用计算里也有奇效。它支持硬件插值、边界处理、缓存优化，特别适合图像处理、查找表这类场景。

特点：

只读
有专门的纹理缓存
支持硬件插值（线性、双线性等）
支持边界处理（clamp、wrap等）
对二维/三维空间局部性友好

我的经验：有一次做图像滤波，用全局内存读像素值，性能一般。换成纹理内存后，因为纹理缓存对二维空间局部性优化得很好，性能提升了30%。而且边界处理都不用自己写代码了，直接设置纹理的寻址模式就行。

// 纹理内存使用示例
texture<float, 2, cudaReadModeElementType> tex;

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

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

3.7 内存模型总结

好了，我把这六种内存整理成一个表格，方便你对比：

内存类型	位置	速度	容量	作用域	生命周期
寄存器	芯片内	最快	极少	单个线程	kernel执行期间
共享内存	芯片内	快	小（KB级）	单个block	block执行期间
本地内存	显存	慢	有限	单个线程	kernel执行期间
全局内存	显存	最慢	大（GB级）	所有线程+CPU	程序运行期间
常量内存	显存（有缓存）	较快（缓存命中时）	小（64KB）	所有线程（只读）	程序运行期间
纹理内存	显存（有缓存）	较快（空间局部性好）	大	所有线程（只读）	程序运行期间

核心原则：尽量把频繁访问的数据放到离计算单元近的地方。寄存器 > 共享内存 > 全局内存。别让数据在慢速内存里反复横跳，那是性能杀手。

这张图把内存层次结构画得很清楚了。从上到下，速度越来越慢，容量越来越大。写kernel的时候，脑子里要有这张图，时刻问自己：我的数据现在在哪？能不能搬到更快的地方去？

好了，这一章的内容就到这里。记住，内存模型是CUDA优化的基石。搞懂了它，后面的算子优化才能事半功倍。

📚 课程目录

第三章：内存模型——全局内存、共享内存、寄存器、本地内存、常量内存、纹理内存

3.1 全局内存（Global Memory）

3.2 共享内存（Shared Memory）

3.3 寄存器（Register）

3.4 本地内存（Local Memory）

3.5 常量内存（Constant Memory）

3.6 纹理内存（Texture Memory）

3.7 内存模型总结