第三章：GPU内存管理——从全局内存到纹理内存的实战指南
说到GPU编程，内存管理绝对是绕不开的核心话题。我刚开始接触CUDA时，总觉得把数据扔到显存里就能跑，结果性能惨不忍睹。后来才明白——不理解内存层次结构，你写的CUDA代码就是碰运气。
这一章，咱们把GPU的几种内存类型掰开揉碎了讲清楚。我会结合自己的踩坑经历，帮你建立正确的内存管理直觉。
3.1 全局内存：最常用，也最容易被滥用
全局内存，说白了就是GPU上的主存。所有线程都能访问，容量最大（通常几个GB到几十GB），但延迟也最高——大约400-800个时钟周期。
我在项目中遇到过一位同事，把临时变量都放在全局内存里，结果kernel跑得比CPU还慢。嗯，这里要注意：全局内存不是用来频繁读写的，它更适合做数据仓库。
核心特点：
- 所有线程可读写
- 生命周期由cudaMalloc/cudaFree控制
- 延迟高，带宽大（但需要合并访问才能发挥）
- 默认的__device__变量存储位置
3.1.1 内存分配与释放
先看最基础的操作——分配和释放。代码很简单，但坑不少。
float *d_data;
size_t size = 1024 * sizeof(float);
// 分配全局内存
cudaMalloc(&d_data, size);
// 使用...（后面会讲拷贝）
// 释放
cudaFree(d_data);我曾经犯过一个低级错误：在循环里反复cudaMalloc/cudaFree，结果程序跑着跑着就崩了。为什么？显存分配是有开销的，频繁分配会产生碎片，最终导致分配失败。
避坑指南：
- 尽量在程序开始一次性分配好所有显存
- 使用cudaMallocPitch处理二维数组（对齐要求）
- 记得检查返回值——cudaError_t不是摆设
- cudaFree(NULL)是安全的，但别依赖这个特性
3.1.2 内存拷贝：cudaMemcpy的三种模式
数据在CPU和GPU之间搬来搬去，全靠cudaMemcpy。你想想看，它其实就三种方向：
| 方向 | 枚举值 | 典型场景 |
|---|---|---|
| 主机→设备 | cudaMemcpyHostToDevice | 把输入数据传到GPU |
| 设备→主机 | cudaMemcpyDeviceToHost | 把计算结果取回CPU |
| 设备→设备 | cudaMemcpyDeviceToDevice | 显存内部数据搬移 |
float *h_data = (float*)malloc(size);
// 初始化h_data...
// 主机到设备
cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);
// 执行kernel...
// 设备到主机
cudaMemcpy(h_data, d_data, size, cudaMemcpyDeviceToHost);这里有个细节：cudaMemcpy是同步操作，它会阻塞CPU直到拷贝完成。我个人习惯用cudaMemcpyAsync配合流（stream）来实现异步传输，后面章节会详细讲。
3.2 共享内存：性能加速的秘密武器
共享内存是GPU上的一块片上存储，延迟只有全局内存的1/10左右（约20-30个时钟周期）。但它有两个限制：
- 容量小：通常每个block只有48KB（部分架构可配到96KB）
- 作用域有限：仅同一个block内的线程可以访问
说白了，共享内存就是线程间的协作空间。我在做矩阵乘法优化时，就是靠共享内存把性能提升了5倍以上。
使用技巧：
- 用__shared__关键字声明
- 记得加__syncthreads()同步——否则数据竞争会让你怀疑人生
- 避免bank conflict（后面会专门讲）
- 动态分配用extern __shared__，大小在kernel启动时指定
__global__ void shared_mem_example(float *input, float *output) {
__shared__ float s_data[256]; // 静态共享内存
int tid = threadIdx.x;
s_data[tid] = input[tid];
__syncthreads(); // 确保所有线程都写完了
// 做一些计算...
output[tid] = s_data[tid] * 2.0f;
}3.3 寄存器：最快但最稀缺
寄存器是GPU上最快的存储，零延迟。但每个线程能用的寄存器数量有限——通常最多255个（取决于计算能力）。
你想想看，如果一个线程用了太多寄存器，编译器会怎么做？它会把多余的变量"溢出"到全局内存，性能直接崩掉。我曾经调试过一个kernel，就因为多声明了几个float变量，性能下降了30%。
寄存器使用建议：
- 用--ptxas-options=-v编译选项查看寄存器使用量
- 尽量减少局部变量数量
- 把循环不变的计算提到循环外
- 考虑用共享内存换寄存器（trade-off）
3.4 常量内存：只读数据的福音
常量内存是全局内存的一个特殊区域，有缓存，适合存储所有线程都访问的只读数据。比如神经网络里的权重参数。
用__constant__声明，通过cudaMemcpyToSymbol拷贝数据。注意：常量内存总大小只有64KB，别想着往里塞大数组。
__constant__ float c_weights[1024];
// 主机端拷贝
float h_weights[1024];
cudaMemcpyToSymbol(c_weights, h_weights, sizeof(h_weights));
// kernel中直接使用
__global__ void use_constant() {
float w = c_weights[threadIdx.x]; // 自动广播到所有线程
}我记得有一次做图像处理，把滤波器的系数放在常量内存里，配合缓存广播机制，性能比全局内存快了将近10倍。
3.5 纹理内存：为空间局部性而生
纹理内存最初是为图形学设计的，但用在通用计算里也有奇效。它有专门的缓存，对二维/三维空间局部性访问特别友好。
使用纹理内存需要：
- 声明纹理引用（texture<float, 2> texRef;）
- 绑定到全局内存（cudaBindTexture）
- 用tex1D/tex2D/tex3D函数读取
- 用完解绑（cudaUnbindTexture）
注意：纹理内存是只读的，而且有最大尺寸限制。现代GPU架构（如Turing、Ampere）更推荐使用只读缓存（__ldg），效果类似但更灵活。
3.6 内存层次结构全景图
说了这么多，咱们用一张图把整个内存体系串起来：
这张图我建议你保存下来。每次写kernel前看一眼，问问自己：我的数据应该放在哪一层？
3.7 实战经验总结
最后，分享几个我这些年积累的实战经验：
- 先规划，后编码：写kernel前，先想清楚哪些数据放共享内存，哪些放寄存器，哪些必须用全局内存。
- 减少全局内存访问：能一次读完的数据，别分多次读。合并访问（coalesced access）是性能关键。
- 共享内存不是万能的：如果数据只用一次，放共享内存反而浪费。只有被多次复用的数据才值得放。
- 用工具说话：nvidia-smi看显存占用，nvcc --ptxas-options=-v看寄存器使用，Nsight Compute看内存带宽利用率。
- 异步拷贝是进阶玩法：cudaMemcpyAsync + 流（stream）可以让数据传输和计算重叠，后面章节会深入讲。
一句话总结：GPU内存管理就是在容量、延迟、作用域之间做权衡。没有银弹，只有根据具体场景选择最合适的存储层次。
嗯，这一章的内容就到这里。记住：内存管理是GPU优化的基石，花时间搞懂它，后面的路会顺畅很多。
公众号:蓝海资料掘金营，微信deep3321