3. 基准测试指标体系与度量方法

做性能基准测试，最容易犯的错误不是“测不准”，而是“测错了东西”。同一段程序，用不同的指标口径去衡量，可能得出完全相反的结论：某个内核的峰值算力很高，但实际任务跑得慢，因为瓶颈在内存带宽；某次测试的吞吐量翻倍，但延迟反而变差，因为批处理把请求攒在了一起。因此，在动手写测试代码之前，必须先建立一套清晰的指标体系，并明确每个指标的测量方法与适用边界。

核心指标的定义与物理含义

FLOPS（每秒浮点运算次数）衡量计算单元的运算吞吐能力，通常以理论峰值和实测值两种形式出现。理论峰值由“每周期浮点运算数 × 核心数 × 频率”估算，例如一个拥有 64 个 FP32 通道、运行在 1.8 GHz 的 GPU 计算单元，其 FP32 理论峰值为 64 × 2 × 1.8×10⁹ ≈ 230 GFLOPS（乘加各计一次运算）。实测 FLOPS 则通过运行高计算强度的内核（如大规模矩阵乘）反推，公式为“总浮点运算量 ÷ 耗时”。两者之间的差距，反映的是流水线利用率、指令调度和寄存器压力等真实约束。

内存带宽描述单位时间内可搬运的数据量，单位 GB/s。它与 FLOPS 共同构成“计算强度”分析的两端：当算法的运算量/字节数低于机器的 FLOPS/带宽比值时，程序就是带宽受限的。测量带宽的经典方法是流式拷贝（读+写）或纯读扫描，用“传输字节数 ÷ 耗时”计算。需要注意，标称带宽（如 HBM 的 1.2 TB/s）是理论值，实测通常只能达到 70%～85%，因为刷新、行切换和访问模式都会带来开销。

延迟是单次操作从发起到完成的绝对时间，单位纳秒或微秒。它关注“一件事多久做完”，与带宽关注的“单位时间做多少事”是不同维度。访存延迟、内核启动延迟、通信延迟都属于此类。延迟对串行依赖强的任务影响巨大，而带宽对大规模数据并行任务更关键。

吞吐量是单位时间内完成的任务数或请求数，单位如 images/s、tokens/s、req/s。它往往由多个资源共同决定，是端到端视角的指标。吞吐量与延迟并非线性关系：提高并发可以提升吞吐量，但可能拉高单请求延迟。

能效比用性能除以功耗表示，如 GFLOPS/W 或 tokens/J。它要求同时采集性能与功耗数据，功耗可通过板卡功耗读数、外接功率计或 CPU 的 RAPL 接口获取。能效比在数据中心场景中越来越重要，但要注意测量窗口应覆盖完整的负载周期，避免只测空闲或峰值瞬间。

测量方法：计时器与性能计数器

最基础的测量工具是计时器。在 CPU 侧，可用 clock_gettime(CLOCK_MONOTONIC, ...) 获取单调时钟，避免系统时间调整带来的跳变；在 GPU 侧，应使用 CUDA Events 或对应的图形 API 查询，而不是简单地在主机端包一个 std::chrono。原因在于 GPU 内核是异步提交的：主机端调用内核后立即返回，若不在流中插入事件并同步，测到的只是“提交耗时”，而非“执行耗时”。正确做法是在内核前后各记录一个事件，调用 cudaEventSynchronize 后计算两者时间差。

cudaEvent_t start, stop;
cudaEventCreate(&start); cudaEventCreate(&stop);
cudaEventRecord(start, stream);
kernel<<<grid, block, 0, stream>>>(d_in, d_out, n);
cudaEventRecord(stop, stream);
cudaEventSynchronize(stop);
float ms = 0; cudaEventElapsedTime(&ms, start, stop);

计时器只能告诉你“花了多久”，性能计数器则能告诉你“时间花在哪”。现代 CPU 提供 PMU（性能监控单元），可读取指令数、周期数、缓存未命中数、分支预测失败数等；GPU 侧有对应的硬件计数器，可通过厂商工具（如 Nsight Compute）采集。用“缓存未命中数 ÷ 总访存次数”可以判断访存模式是否友好，用“实际指令数 ÷ 理论指令数”可以评估流水线效率。性能计数器的价值在于把黑盒的耗时拆解为可归因的硬件事件。

常见陷阱与规避方法

异步执行陷阱是最常见的问题。除了上述 GPU 内核异步，还有 DMA 拷贝、多流并发等场景。规避原则是：测量区间必须包含完整的同步点，且同步点应放在被测操作之后、计时结束之前。若测试多流并发，需确保所有流都完成后再停止计时。

缓存效应会系统性地高估性能。若测试数据规模小于末级缓存，重复访问会命中缓存，测出的带宽可能是真实内存带宽的数倍。规避方法是让工作集显著大于缓存容量（例如用数百 MB 到数 GB 的数组），并在多轮迭代中轮换数据地址，避免固定缓冲区被反复命中。同时要区分“冷启动”和“热态”两种口径，分别报告。

预热不足同样会造成偏差。频率提升、TLB 填充、JIT 编译、显存页表建立都需要时间，第一轮往往偏慢。通常应丢弃前若干轮，取稳定后的中位数或最小值。此外，计时粒度不宜过粗：若单次操作只有几微秒，而计时器分辨率是微秒级，相对误差会很大，应通过批量重复取平均来摊薄误差。

一个可复用的度量记录表

为了让每次测试都可追溯，建议为每个基准用例记录如下字段。下表为示例结构，数值为合理假设，仅用于说明填写方式：

指标体系的真正作用，是让“快”和“慢”变成可解释的判断。当实测 FLOPS 远低于峰值时，去看计算强度是否偏低、是否带宽受限；当带宽利用率不足时，去查访问是否连续、缓存是否被浪费；当吞吐量提升而延迟恶化时，去确认并发策略是否以牺牲响应性为代价。把指标、测量方法和陷阱三者对应起来，基准测试才具备工程决策价值，而不只是一串好看的数字。