第三章：CUDA内存优化——全局内存合并访问、共享内存bank conflict、常量内存与纹理内存、统一内存使用技巧
内存优化，说白了就是让数据在GPU里跑得顺畅。我见过太多人把精力全放在计算上，结果性能瓶颈卡在内存访问上。嗯，今天咱们就把这块硬骨头啃下来。
3.1 全局内存合并访问
全局内存是GPU里最大的存储空间，但也是最慢的。为什么慢？因为它离计算单元远。我个人习惯把全局内存想象成一个巨大的仓库，每次取货都要花时间。
合并访问的核心思想：让同一个warp里的32个线程，访问连续的内存地址。这样硬件可以把多次访问合并成一次大的传输。
关键点：合并访问要求对齐到128字节。说白了，就是起始地址得是128的倍数。
来看个例子。假设我们要处理一个float数组：
// ❌ 糟糕的访问模式：按列访问
for (int i = 0; i < N; i++) {
float val = data[threadIdx.x + i * blockDim.x]; // 跨步访问
}
// ✅ 好的访问模式：按行访问
for (int i = 0; i < N; i++) {
float val = data[i + threadIdx.x * N]; // 连续访问
}我在项目中遇到过一位同事，他写了个矩阵转置，性能惨不忍睹。一查，就是按列访问全局内存。改成按行访问后，速度提升了将近10倍。你想想看，就改了个索引顺序，效果天差地别。
小技巧：使用cudaMemcpy时，尽量保证数据在主机端和设备端都是对齐的。我一般用cudaMallocPitch来分配二维数组，它能自动处理对齐。
3.2 共享内存bank conflict
共享内存是片上的，速度比全局内存快两个数量级。但它有个坑——bank conflict。
共享内存被分成32个bank，每个bank宽度4字节。如果同一个warp里的多个线程访问同一个bank的不同地址，就会发生冲突。冲突会导致访问串行化，性能下降。
为什么会这样？ 硬件设计上，每个bank每个周期只能服务一个请求。如果两个线程撞到同一个bank，就得排队。
| 访问模式 | 是否冲突 | 性能影响 |
|---|---|---|
| 所有线程访问不同bank | 无冲突 | 最佳 |
| 多个线程访问同一bank的不同地址 | 冲突 | 串行化 |
| 所有线程访问同一bank的同一地址 | 广播 | 无冲突（硬件支持广播） |
我曾经调试过一个卷积核，共享内存访问模式没设计好，导致bank conflict严重。后来我用了padding技巧——在共享内存数组的每一行末尾加几个空元素，把数据错开。就这么简单的一招，性能提升了30%。
// 避免bank conflict的padding技巧
__shared__ float s_data[32][32 + 1]; // 每行多一个元素注意：padding会增加共享内存使用量，别加太多。我一般每行加1-4个元素就够了。
3.3 常量内存与纹理内存
这两个是GPU里的特殊存储区域，用好了能省不少事。
常量内存：容量小（64KB），但有专门的缓存。适合存储不会变的数据，比如滤波器系数、查找表。所有线程访问同一个地址时，常量内存最快。
我个人习惯把模型参数放在常量内存里。记得有一次做图像处理，把3x3的卷积核放常量内存，比放全局内存快了将近2倍。
__constant__ float kernel[9]; // 声明常量内存
// 主机端复制
cudaMemcpyToSymbol(kernel, host_kernel, 9 * sizeof(float));纹理内存：有硬件插值功能，适合图像处理。它支持二维和三维的局部性访问，而且对非对齐访问很友好。
嗯，这里要注意：纹理内存的延迟比共享内存高，但比全局内存低。如果你做的是图像滤波、插值这类操作，纹理内存是首选。
避坑指南：我曾经在纹理内存上吃过亏——忘记设置归一化坐标。结果纹理坐标范围搞错了，图像显示全乱套。记得检查texture对象的参数设置。
3.4 统一内存使用技巧
统一内存（Unified Memory）是CUDA 6.0引入的，它让CPU和GPU共享同一个地址空间。说白了，你不需要手动管理数据迁移了。
但别高兴太早。统一内存不是银弹，用不好反而更慢。
什么时候用统一内存？
- 代码原型开发阶段，快速验证想法
- 数据访问模式复杂，手动管理内存太麻烦
- 处理不规则数据结构（比如链表、树）
什么时候别用？
- 性能敏感的核心循环
- 数据访问模式很规律（这时候手动管理更快）
- 内存容量紧张（统一内存有额外开销）
// 统一内存分配
float *data;
cudaMallocManaged(&data, N * sizeof(float));
// CPU和GPU都能直接访问
// 系统会自动迁移数据我有个项目，数据是稀疏图结构，用统一内存写起来特别爽。但后来发现性能瓶颈在页面迁移上。解决方案是：用cudaMemPrefetchAsync提前把数据迁移到GPU，避免运行时缺页中断。
核心技巧：使用cudaMemAdvise设置内存访问建议，告诉系统哪些数据会被频繁访问。这能减少不必要的页面迁移。
你想想看，统一内存就像自动挡汽车——市区开很方便，但赛道日还是手动挡更快。选哪个，取决于你的场景。
好了，这一章的内容就这些。记住：内存优化的核心是理解硬件的访问模式，然后让你的代码去适应它。别跟硬件较劲，顺着它来，性能自然就上去了。