3、CUDA内存优化：内存合并访问、共享内存的使用与bank conflict、常量内存与纹理内存、统一内存（Unified Memory）的使用场景

各位同学，咱们今天聊点实在的。GPU 加速，说白了就是跟内存打交道。你 kernel 写得再花哨，访存模式一塌糊涂，性能立马打回原形。我这些年调优踩过的坑，十有八九都跟内存有关。今天就把这几个核心点掰开揉碎讲清楚。

3.1 内存合并访问：让带宽跑满的秘诀

先问个问题：为什么同样的计算量，有的人 kernel 跑得飞快，有的人却慢得像蜗牛？

答案往往在访存模式上。GPU 的全局内存带宽很大，但有个前提——你得让同一 warp 里的 32 个线程，访问连续的地址空间。这就是所谓的「内存合并访问」。

我举个例子。假设你要读一个 float 数组：

// 好的写法：合并访问
float val = data[threadIdx.x];  // 线程0读data[0]，线程1读data[1]...

// 坏的写法：非合并访问
float val = data[threadIdx.x * 64];  // 线程0读data[0]，线程1读data[64]...

第一种写法，32 个线程访问 32 个连续的 float，正好 128 字节。硬件一次事务就能搞定。

第二种写法，跨度 64 个 float，也就是 256 字节。硬件得发起多次事务，带宽利用率直线下降。

核心原则：让 warp 内相邻线程访问相邻地址。

我在项目中遇到过一个问题：一个矩阵转置的 kernel，性能死活上不去。后来一查，发现列方向读取时，线程 stride 等于矩阵行宽，完全没合并。改成共享内存做分块转置后，性能直接翻了 4 倍。

避坑指南：二维数组的访问顺序。C/C++ 是行优先，所以最内层循环应该遍历列方向。如果你写成列优先访问，那就等着性能崩盘吧。

3.2 共享内存：片上高速缓存的艺术

共享内存是 GPU 上的 SRAM，延迟比全局内存低一个数量级。但用好它，需要点技巧。

共享内存的典型用法是做数据复用。比如矩阵乘法，每个线程块把子矩阵加载到共享内存里，然后反复使用，避免反复访问全局内存。

__shared__ float tile_A[BLOCK_SIZE][BLOCK_SIZE];
__shared__ float tile_B[BLOCK_SIZE][BLOCK_SIZE];

// 协作加载
tile_A[threadIdx.y][threadIdx.x] = A[row * W + tx];
tile_B[threadIdx.y][threadIdx.x] = B[ty * W + col];
__syncthreads();

// 从共享内存计算
for (int k = 0; k < BLOCK_SIZE; ++k) {
    sum += tile_A[threadIdx.y][k] * tile_B[k][threadIdx.x];
}

这里有个关键点：__syncthreads() 不能少。它保证所有线程都加载完了，才能开始计算。我见过有人漏掉这个同步，结果数据还没到位就开始算，出来的结果全是错的。

3.2.1 Bank Conflict：共享内存的隐形杀手

共享内存被分成 32 个 bank（对于计算能力 6.x 以上的设备）。每个 bank 每周期只能响应一个地址。如果同一 warp 里多个线程访问同一个 bank 的不同地址，就会发生 bank conflict，导致访问串行化。

举个例子：

// 有 bank conflict：线程0访问[0]，线程1访问[32]（都在bank 0）
float val = shared[threadIdx.x * 32];

// 无 bank conflict：线程0访问[0]，线程1访问[1]（不同bank）
float val = shared[threadIdx.x];

小技巧：如果访问步长是奇数，通常不会产生 bank conflict。因为 32 个 bank 的映射是地址对 32 取模，奇数步长会让线程均匀分布在各个 bank 上。

我曾经在实现一个 3D 卷积算子时，共享内存的 bank conflict 导致性能损失了 30%。后来把数据布局从 [z][y][x] 改成 [x][y][z] 的 padding 方式，才把冲突消除掉。嗯，这种细节调试起来确实挺磨人的。

3.3 常量内存与纹理内存：各有各的用武之地

常量内存和纹理内存，很多人觉得鸡肋。其实用对场景，效果出奇的好。

3.3.1 常量内存

常量内存只有 64KB，但有两个优势：

  • 有专用的常量缓存，如果同一 warp 所有线程访问同一个地址，延迟极低（一个周期）
  • 数据是只读的，编译器可以做更多优化

适合场景：查找表、滤波器系数、模型参数等所有线程都读相同值的情况。

__constant__ float filter[256];

// 主机端拷贝
cudaMemcpyToSymbol(filter, host_filter, sizeof(float) * 256);

// 设备端使用：所有线程读同一个 filter[idx]
float val = data[i] * filter[idx];

我建议：如果你的 kernel 里有大量只读的、所有线程共享的常量数据，优先考虑常量内存。64KB 虽然小，但足够放很多模型参数了。

3.3.2 纹理内存

纹理内存最初是为图形学设计的，但在深度学习里也有妙用。它支持：

  • 硬件插值（对图像缩放有用）
  • 边界处理（clamp、wrap 等模式）
  • 对二维/三维空间局部性好的访问模式有缓存优化

说实话，在纯计算场景里，纹理内存用得不多。但如果你做图像预处理、特征图的上采样，纹理内存的硬件插值能省不少事。

注意：纹理内存的读取延迟比全局内存高，但缓存命中率高。如果你的访问模式是随机的，纹理内存反而可能更慢。

3.4 统一内存（Unified Memory）：简化编程的双刃剑

统一内存（UM）是 CUDA 6.0 引入的特性。它让 CPU 和 GPU 共享同一个虚拟地址空间。你不需要手动 cudaMemcpy，系统会自动迁移数据。

// 统一内存：分配一次，两边都能访问
float *data;
cudaMallocManaged(&data, N * sizeof(float));

// CPU 端初始化
for (int i = 0; i < N; i++) data[i] = i;

// GPU 端使用
kernel<<<grid, block>>>(data, N);

cudaDeviceSynchronize();  // 自动迁移完成

听起来很美好对吧？但我要泼点冷水。

使用场景一：原型开发

我刚开始用 UM 时觉得太爽了，再也不用管内存拷贝了。但后来发现，性能上不去。为什么？因为 UM 的数据迁移是 page fault 驱动的。你访问一个不在当前设备上的页面，就会触发缺页中断，然后数据才迁移过来。这个过程很慢。

使用场景二：不规则数据结构

比如链表、树、图这种指针满天飞的数据结构。用传统方式，你得手动把整个结构体序列化再拷贝，非常痛苦。UM 直接搞定，指针在两边都有效。

使用场景三：数据访问模式稀疏且不可预测

如果你的程序访问数据的位置完全随机，手动管理内存拷贝反而更难优化。UM 的按需迁移反而可能更高效。

我的建议：能用显式拷贝就用显式拷贝。UM 只适合原型验证、不规则数据结构、或者数据量小到可以忽略迁移开销的场景。

我曾经在一个推荐系统项目里，图数据用 UM 管理。数据量 2GB，每次训练都要触发大量 page fault，性能比手动管理慢了 5 倍。后来改成手动分块拷贝，才把性能拉回来。说白了，UM 是方便，但性能上你得有心理准备。

小结

今天聊了四个内存优化点：

  • 合并访问：让 warp 内线程访问连续地址，这是性能的基石
  • 共享内存：片上高速缓存，注意 bank conflict 和同步
  • 常量/纹理内存：特定场景下的利器，别滥用
  • 统一内存：方便但有代价，选对场景再用

你想想看，这些知识点单独拎出来都不难，但组合在一起，就能决定你的 kernel 是跑 100μs 还是 500μs。下次写 kernel 时，先问问自己：我的访存模式对吗？共享内存用好了吗？有没有更合适的内存类型？

嗯，今天就到这里。下一章咱们聊聊算子融合和 kernel 合并，那又是另一番天地了。

← 上一章 📖 返回目录 下一章 →