4. 性能指标：延迟、带宽、消息速率、吞吐量的定义与测量方法

做NVSHMEM性能调优，说白了就是跟四个数字打交道：延迟、带宽、消息速率、吞吐量。我刚开始接触这玩意儿的时候，也搞混过它们之间的关系。今天咱们就把这四个指标掰开揉碎了讲清楚。

4.1 延迟（Latency）—— 消息的“响应速度”

定义：从一个PE发起操作，到目标PE收到数据并返回确认，所经历的时间。单位通常是微秒（μs）或纳秒（ns）。

我习惯把延迟比作“快递小哥从你家到邻居家的往返时间”。你发一个包裹，邻居签收后告诉你收到了，这个总耗时就是延迟。

关键点：延迟是单次消息的“反应时间”，跟消息大小关系不大。哪怕你只发1个字节，该走的握手流程一个都不会少。

测量方法：

  • 使用NVSHMEM的nvshmem_barrier_all()做同步
  • 用nvshmem_double_wait_until()配合高精度计时器
  • 经典做法：ping-pong测试，记录往返时间除以2
// 延迟测量示例：ping-pong
double start, end;
int peer = (my_pe + 1) % n_pes;

for (int i = 0; i < WARMUP; i++) {
    nvshmem_double_p(&data, 1.0, peer);
    nvshmem_double_wait_until(&data, NVSHMEM_CMP_NE, 0.0);
}

start = MPI_Wtime();
for (int i = 0; i < ITERATIONS; i++) {
    nvshmem_double_p(&data, 1.0, peer);
    nvshmem_double_wait_until(&data, NVSHMEM_CMP_NE, 0.0);
}
end = MPI_Wtime();

double latency = (end - start) / ITERATIONS / 2.0;  // 除以2得到单程延迟
printf("Latency: %.2f us\n", latency * 1e6);

避坑指南：我曾经在测量延迟时忘了做WARMUP，结果前几次调用包含了驱动加载和内存注册的开销，数据直接偏大了30%。记住，一定要先跑几轮“热身”。

4.2 带宽（Bandwidth）—— 管道的“粗细”

定义：单位时间内能传输的数据量。单位是GB/s或MB/s。

你想想看，延迟是“快递小哥跑一趟的时间”，带宽就是“这条路上能并排跑多少辆快递车”。带宽越大，同样时间内能搬的数据就越多。

关键点：带宽跟消息大小强相关。小消息（比如8字节）带宽上不去，因为延迟开销占大头。大消息（比如1MB）才能跑满带宽。

测量方法：

  • 使用nvshmem_putmem()或nvshmem_getmem()传输大块数据
  • 记录传输时间，用数据量除以时间
  • 注意区分单向带宽和双向带宽
// 带宽测量示例：单向put
size_t msg_size = 1 << 20;  // 1MB
char *src = nvshmem_malloc(msg_size);
char *dst = nvshmem_malloc(msg_size);

start = MPI_Wtime();
for (int i = 0; i < ITERATIONS; i++) {
    nvshmem_putmem(dst, src, msg_size, peer);
    nvshmem_quiet();  // 确保完成
}
end = MPI_Wtime();

double bw = (double)msg_size * ITERATIONS / (end - start);
printf("Bandwidth: %.2f GB/s\n", bw / 1e9);

注意：测量带宽时，nvshmem_quiet()是必须的。我见过有人忘了加，结果测出来的带宽虚高，因为put操作还没真正完成就计时结束了。

4.3 消息速率（Message Rate）—— 小消息的“爆发力”

定义：每秒能处理的消息数量。单位是Mpps（百万消息每秒）。

说白了，消息速率就是“快递小哥每秒能跑多少趟”。它跟延迟互为倒数关系，但又不完全一样。延迟是单趟时间，消息速率是每秒趟数。

关键点：消息速率主要受限于硬件处理能力，比如NVLink的包处理引擎、PCIe的TLP包数量。小消息（8-64字节）最能体现消息速率。

测量方法：

  • 使用nvshmem_short_put()或nvshmem_double_p()发小消息
  • 统计单位时间内完成的操作次数
  • 注意不要用nvshmem_quiet()，否则会串行化
// 消息速率测量示例
start = MPI_Wtime();
for (int i = 0; i < ITERATIONS; i++) {
    nvshmem_int_p(&data, i, peer);
    // 注意：这里不用quiet，让操作流水线化
}
// 最后统一等待完成
nvshmem_quiet();
end = MPI_Wtime();

double msg_rate = ITERATIONS / (end - start);
printf("Message Rate: %.2f Mpps\n", msg_rate / 1e6);

个人经验：我曾经在A100集群上测过，NVSHMEM的消息速率能到40-50 Mpps。但如果你用了nvshmem_quiet()，速率直接掉到个位数。原因很简单——每次等待都打断了流水线。

4.4 吞吐量（Throughput）—— 综合的“搬运能力”

定义：单位时间内完成的有效工作量。对于NVSHMEM来说，就是每秒传输的有效数据量。单位是GB/s。

嗯，这里要注意区分吞吐量和带宽。带宽是“管道能跑多快”，吞吐量是“实际跑了多少货”。比如你带宽有100GB/s，但应用层协议开销大，实际吞吐量可能只有60GB/s。

关键点：吞吐量是综合指标，它包含了延迟、带宽、消息速率的影响。实际应用中，我们最关心的就是吞吐量。

测量方法：

  • 模拟真实应用场景，比如Stencil计算、Alltoall通信
  • 记录总传输数据量和总耗时
  • 吞吐量 = 有效数据量 / 总时间
// 吞吐量测量示例：模拟Alltoall
size_t chunk_size = 1 << 12;  // 4KB per pair
size_t total_bytes = chunk_size * n_pes * (n_pes - 1);

start = MPI_Wtime();
for (int i = 0; i < n_pes; i++) {
    if (i == my_pe) continue;
    nvshmem_putmem(dst[i], src[i], chunk_size, i);
}
nvshmem_barrier_all();  // 等待所有完成
end = MPI_Wtime();

double throughput = total_bytes / (end - start);
printf("Throughput: %.2f GB/s\n", throughput / 1e9);

避坑指南：我曾经在测吞吐量时，用了nvshmem_quiet()而不是nvshmem_barrier_all()。结果发现吞吐量只有预期的60%。后来才意识到，quiet()只保证本地操作完成，不保证远端也完成了。对于吞吐量这种需要全局同步的场景，一定要用barrier。

4.5 四个指标的关系与权衡

这四个指标不是孤立的。我画了张图，帮你理清它们之间的关系：

延迟 Latency 带宽 Bandwidth 消息速率 Message Rate 吞吐量 Throughput 倒数关系 上限约束 小消息贡献 基础开销 图：NVSHMEM四大性能指标关系图

从图中你能看出：

  • 延迟是基础，它决定了消息速率的上限
  • 带宽是大消息场景下的天花板
  • 消息速率在小消息场景下起主导作用
  • 吞吐量是前三者的综合体现

实际调优建议：

  • 小消息（< 256字节）：关注延迟和消息速率
  • 中等消息（256字节 - 64KB）：平衡考虑
  • 大消息（> 64KB）：关注带宽

4.6 测量时的常见陷阱

我踩过的坑不少，挑几个典型的说说：

  1. 计时器精度不够：用clock_gettime()代替gettimeofday()，后者精度只有微秒级，对于纳秒级的延迟测量不够用。
  2. 忘记同步：测量延迟时，远端还没收到数据你就开始计时了，结果偏小。测量带宽时，数据还没传完你就停了，结果偏大。
  3. 缓存污染：连续测量时，数据可能还在L1/L2缓存里。我习惯在每次测量前清空缓存，或者用随机地址访问。
  4. PCIe链路竞争：多GPU共享PCIe Switch时，带宽会被分摊。测量时要确认链路独占。

我的测量清单：每次做基准测试前，我都会检查这几点——GPU时钟是否锁定、PCIe链路是否Gen4 x16、NVLink是否全部激活、有没有其他进程抢占GPU。这些细节不注意，测出来的数据就是废的。

好了，四个指标的定义和测量方法就讲到这里。记住一句话：延迟看小消息，带宽看大消息，消息速率看爆发力，吞吐量看综合实力。实际调优时，根据你的应用场景，找到最关键的指标去优化就行。

一句话总结：延迟是“快不快”，带宽是“多不多”，消息速率是“频不频”，吞吐量是“实不实”。

公众号:蓝海资料掘金营，微信deep3321

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