4、CUDA内存优化：全局内存合并访问、共享内存bank conflict、内存带宽利用率分析

说到CUDA性能优化，内存这块绝对是绕不开的坎。我见过太多同学，核函数写得挺漂亮，一跑起来带宽利用率不到20%。说白了，GPU算力再强，喂不饱数据也是白搭。今天咱们就聊聊内存优化的三个核心点：全局内存的合并访问、共享内存的bank conflict，还有带宽利用率怎么分析。

4.1 全局内存合并访问

先问个问题：你知道GPU从全局内存读数据，一次能读多少吗？答案是32字节、64字节或128字节，取决于你的访问模式。嗯，这里要注意——如果同一个warp里的32个线程，访问的地址是连续的，那硬件就能把这些请求合并成一次大的内存事务。这就是所谓的「合并访问」。

我举个例子。假设你有一个float数组，每个线程要读一个float（4字节）。如果线程0读data[0]，线程1读data[1]……线程31读data[31]，那这32个地址正好是连续的128字节。硬件一次事务就搞定了，完美！

反过来，如果线程0读data[0]，线程1读data[32]，线程2读data[64]……那每个线程的地址间隔32个float，128字节。硬件就得拆成32次事务，效率直接掉到3%左右。我在项目中遇到过这种坑，当时一个简单的矩阵转置，带宽利用率只有5%，改成合并访问后直接飙到70%。

// ❌ 非合并访问（按列访问）
__global__ void bad_access(float *A, float *B, int N) {
    int row = blockIdx.x * blockDim.x + threadIdx.x;
    int col = blockIdx.y * blockDim.y + threadIdx.y;
    if (row < N && col < N) {
        B[col * N + row] = A[row * N + col];  // 注意这里：col * N + row
    }
}

// ✅ 合并访问（按行访问）
__global__ void good_access(float *A, float *B, int N) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < N && col < N) {
        B[row * N + col] = A[row * N + col];  // 连续地址
    }
}

4.2 共享内存bank conflict

共享内存是片上的，速度比全局内存快两个数量级。但它有个脾气——bank conflict。共享内存被分成32个bank（对应warp的32个线程），每个bank宽度4字节。如果同一个warp里多个线程访问同一个bank的不同地址，就会发生冲突，访问会被串行化。

你想想看，本来一次就能读完的数据，现在要拆成好几次，性能自然就下来了。那什么情况下不会冲突呢？

• 无冲突：每个线程访问不同的bank，或者所有线程访问同一个地址（广播）。
• 2路冲突：两个线程访问同一个bank的不同地址，性能减半。
• n路冲突：n个线程访问同一个bank的不同地址，性能降到1/n。

我记得有一次做矩阵乘法优化，用共享内存存分块数据。一开始直接按行存，结果在读取的时候发现bank conflict严重。后来我用了padding技巧——每行末尾加一个float的填充，把bank对齐打散，冲突就没了。

// 共享内存声明，加padding避免bank conflict
__shared__ float s_data[BLOCK_SIZE][BLOCK_SIZE + 1];  // 注意+1

// 写入时
s_data[threadIdx.y][threadIdx.x] = A[row * N + col];

// 读取时，每个线程访问不同的列，自然就分散到不同bank了

4.3 内存带宽利用率分析

带宽利用率，说白了就是你的核函数实际用到了多少理论带宽。计算公式很简单：

利用率 = 实际带宽 / 理论带宽 × 100%

实际带宽怎么算？用NVIDIA的profiler工具就行。我个人习惯用ncu（Nsight Compute），它能直接告诉你Memory Throughput和Memory Utilization。

那什么样的利用率算好呢？我一般这么看：

利用率范围	评价	建议
> 80%	优秀	基本到顶了，别折腾了
50% - 80%	良好	还有优化空间，检查访问模式
20% - 50%	一般	大概率有合并访问问题或bank conflict
< 20%	糟糕	赶紧检查代码，可能访问模式全错了

为什么会利用率低？我总结了几种常见原因：

1. 非合并访问：warp内线程地址不连续，这是最常见的。
2. bank conflict：共享内存访问冲突，导致有效带宽下降。
3. 计算密度太低：每个数据只做一次运算就写回，带宽被浪费在搬运上。
4. 线程束发散：warp内线程走不同分支，部分线程闲置。

4.4 实战：用profiler定位问题

光说不练假把式。咱们用ncu跑一下前面那个非合并访问的例子，看看数据怎么说。

# 编译
nvcc -o bad_access bad_access.cu

# 用ncu分析
ncu --set full -o bad_access_report ./bad_access

# 关键输出指标
# Memory Throughput: 12.3 GB/s (理论带宽900 GB/s，利用率1.4%)
# L1/TEX Cache Throughput: 8.7 GB/s
# Sector Misses: 98.7%  ← 这个指标说明几乎每次访问都没命中

看到没？Sector Misses 98.7%，说明几乎每次内存访问都没能合并。改成合并访问后：

# 优化后
# Memory Throughput: 720.5 GB/s (利用率80.1%)
# Sector Misses: 2.3%

从1.4%到80.1%，就改了一行代码。嗯，这就是内存优化的魅力。

好了，这一章的内容就到这儿。下一章咱们聊聊「CUDA流与并发执行」，到时候我会分享一个用多流隐藏延迟的实战案例，挺有意思的。