3、内存带宽理论极限：如何计算GPU的理论峰值带宽，为什么实际跑不满？

好，咱们来聊聊一个很实在的问题。

你写了个CUDA程序，数据拷贝花了半天。你心想，这GPU带宽不是标称好几百GB/s吗？怎么我测出来就这点？

嗯，这里有个坑。你看到的那个数字，叫理论峰值带宽。你实际能拿到的，叫有效带宽。两者之间，隔着一条鸿沟。

我个人习惯，拿到一块新GPU，第一件事就是算算它的理论带宽。不是为了背参数，而是为了心里有个底——我到底离天花板还有多远？

3.1 理论峰值带宽的计算公式

公式其实很简单，就一个乘法：

理论带宽 = 显存时钟频率 × 数据总线位宽 × 传输倍率

拆开来看：

• 显存时钟频率：单位是MHz或GHz。比如GDDR6常见的1750 MHz。
• 数据总线位宽：单位是bit。比如RTX 3090是384-bit。
• 传输倍率：这是关键。DDR内存是2倍，GDDR6是4倍，GDDR6X是8倍（甚至更高）。

举个例子。我手头有块RTX 3090，参数如下：

• 显存频率：1950 MHz
• 位宽：384-bit
• GDDR6X，倍率8

计算过程：

1950 MHz × 384 bit × 8 ÷ 8 (bit转Byte) ÷ 1000 (MHz转GHz) = 748.8 GB/s

嗯，跟官方标称的936 GB/s差了点？别急，官方用的是有效频率。GDDR6X的有效频率是19.5 Gbps，直接乘位宽再除8就行：

19.5 Gbps × 384 bit ÷ 8 = 936 GB/s

你看，对上号了。

3.2 为什么实际跑不满？——三大瓶颈

理论值算出来了，但实际跑分往往只有理论值的60%-80%。为什么会这样？

我在项目中遇到过好几次，明明带宽算得挺高，一测就傻眼。后来总结出三个主要原因：

3.2.1 协议开销

说白了，数据传输不是光着膀子跑的。它要穿衣服——协议头、校验位、地址信息等等。这些额外数据占用了宝贵的带宽。

• PCIe协议：每次传输有TLP头、DLLP头，大概占5%-10%的开销。
• 显存控制器：行激活、列选通、预充电，这些操作都在消耗时钟周期。

你想想看，你付了100%的带宽钱，结果有10%在干杂活。

3.2.2 访存模式不理想

这是大头。GPU喜欢连续、对齐的访问。你如果东一榔头西一棒子，带宽直接腰斩。

• 合并访问：相邻线程访问相邻地址，带宽利用率最高。
• 随机访问：每个线程访问不同page，显存控制器要反复换行，效率暴跌。

我曾经调试过一个程序，带宽只有理论值的20%。查了半天，发现是索引数组没排序，导致全局内存访问完全随机。改成排序后，带宽直接飙到75%。

3.2.3 硬件限制

嗯，这里要注意。理论带宽是理想情况下的极限。实际上，显存控制器、片上互联、L2缓存，都有各自的瓶颈。

• 显存控制器数量：比如A100有40个控制器，但每个控制器有自己的频率上限。
• 跨片访问：多GPU或MCM（多芯片模块）架构，跨片访问带宽远低于片内。
• 温度降频：跑久了温度上来，显存频率自动下降，带宽也跟着掉。

3.3 如何测量实际带宽？

别光算，动手测一测。CUDA自带的bandwidthTest工具就很好用。

# 编译
cd /usr/local/cuda/samples/1_Utilities/bandwidthTest
make

# 运行
./bandwidthTest

# 输出示例
[CUDA Bandwidth Test] - Starting...
Running on...
Device 0: Tesla V100-SXM2-32GB
Quick Mode
Host to Device Bandwidth: 11.76 GB/s
Device to Host Bandwidth: 12.55 GB/s
Device to Device Bandwidth: 836.22 GB/s

你看，V100的理论带宽是900 GB/s，实测836 GB/s，利用率约93%。这已经非常高了，因为bandwidthTest用的是最优的连续拷贝模式。

但换成你自己的程序，可能就惨不忍睹了。我建议你写个简单的kernel，测测不同访存模式下的带宽：

// 连续访问（高带宽）
__global__ void copy_contiguous(float *in, float *out, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) out[idx] = in[idx];
}

// 跨步访问（低带宽）
__global__ void copy_stride(float *in, float *out, int N, int stride) {
    int idx = (blockIdx.x * blockDim.x + threadIdx.x) * stride;
    if (idx < N) out[idx] = in[idx];
}

跑一下，你会发现跨步访问的带宽可能只有连续访问的十分之一。

3.4 一张图看懂带宽瓶颈

下面这张SVG图，是我自己画的一个带宽分析框架。它把从理论值到实际值的损耗过程，拆成了几个关键环节。

3.5 小结：别被理论值骗了

理论带宽是个参考，不是目标。我见过太多人，拿着理论值去评估程序性能，结果发现差了一大截，就开始怀疑人生。

其实，你只要做到以下几点，就能拿到不错的实际带宽：

• 用cudaMalloc分配内存，保证对齐。
• 尽量连续访问，避免跨步和随机。
• 一次拷贝大块数据，别小片小片地拷。
• 用cudaMemcpyAsync，配合流和事件，让拷贝和计算重叠。

嗯，做到这些，你的带宽利用率大概率能到70%以上。剩下的，就是硬件和协议的硬限制了，别太纠结。