4. 吞吐压测方法论

做GPU性能优化这些年，我踩过最大的坑就是——测试结果不准。你辛辛苦苦改了一版kernel，跑出来快了20%，结果换个参数又慢了。说白了，问题出在测试方法上。

吞吐压测不是随便跑几个时间就完事的。它需要一套严谨的方法论。我管它叫「微基准测试设计原则」。今天咱们就聊聊这个。

为什么要做微基准测试？

你想想看，一个完整的AI训练脚本，里面几十个算子。你改了一个kernel，整体时间变了5%。这5%是来自你的优化，还是别的因素？很难说清楚。

微基准测试就是把单个kernel拎出来，单独测它的吞吐。这样你才能精准判断：这个优化到底有没有用。

我在项目中遇到过好几次，整体跑得快了，但其实是数据加载那边变快了，跟我的kernel优化半毛钱关系没有。嗯，从那以后，我坚持每个kernel单独做微基准测试。

消除启动开销：warm-up

GPU第一次启动kernel时，有各种初始化开销。包括驱动加载、上下文创建、内存分配等等。这些开销跟你的kernel性能无关，但会污染测试结果。

怎么处理？做warm-up。

我一般这样写测试代码：

// warm-up: 跑5次，不计时
for (int i = 0; i < 5; i++) {
    myKernel<<<grid, block>>>(d_data);
}
cudaDeviceSynchronize();

// 正式计时：跑100次取中位数
float times[100];
for (int i = 0; i < 100; i++) {
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    
    cudaEventRecord(start);
    myKernel<<<grid, block>>>(d_data);
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    
    cudaEventElapsedTime(&times[i], start, stop);
    
    cudaEventDestroy(start);
    cudaEventDestroy(stop);
}

warm-up的次数，我个人习惯设5到10次。太少不够热，太多浪费时间。你根据kernel的复杂度调整就好。

多次运行取中位数

为什么取中位数，不取平均值？

因为GPU运行时间有波动。偶尔一次调度延迟，或者系统中断，会导致某个时间点特别长。平均值会被这些异常值拉高，中位数不会。

举个例子：

运行次数	时间(ms)
1-9次	1.02, 1.01, 1.03, 1.02, 1.01, 1.02, 1.03, 1.01, 1.02
第10次	2.15（被调度打断了）

平均值是1.13ms，中位数是1.02ms。哪个更接近真实性能？显然是1.02ms。

改变grid/block大小

同一个kernel，用不同的grid和block大小，吞吐可能差好几倍。你想想看，这能叫「稳定」的测试吗？

所以，做吞吐压测时，要扫一遍grid/block参数空间。我一般这样设计：

// 测试不同的block大小
int block_sizes[] = {64, 128, 256, 512, 1024};
int grid_sizes[] = {1, 2, 4, 8, 16, 32, 64, 128};

for (int b = 0; b < 5; b++) {
    for (int g = 0; g < 8; g++) {
        dim3 block(block_sizes[b]);
        dim3 grid(grid_sizes[g]);
        
        // 执行测试并记录吞吐
        float throughput = testKernel(grid, block);
        printf("grid=%d, block=%d, throughput=%.2f GB/s\n", 
               grid_sizes[g], block_sizes[b], throughput);
    }
}

然后画一张热力图，横轴是grid大小，纵轴是block大小，颜色代表吞吐。一眼就能看出最佳配置在哪。

控制变量法测试

这是最基础，但也最容易忽略的原则。每次只改一个变量，其他全部固定。

比如你想测「共享内存对性能的影响」：

• 固定grid/block大小
• 固定数据类型
• 固定数据量
• 只改变是否使用共享内存

这样你才能说：性能差异是共享内存带来的。

我见过有人同时改了block大小和数据类型，然后说「优化后快了30%」。你问他到底是哪个优化起了作用？他说不清楚。这其实是不专业的。

知识体系总览

下面这张图，是我对吞吐压测方法论的总结。你可以把它当作一个checklist，每次做测试前过一遍：

避坑指南

最后，分享几个我踩过的坑：

• 不要用第一次运行的结果——我刚开始做测试时，直接拿第一次运行的时间当结果。后来发现每次重启程序后第一次运行都特别慢。这就是没做warm-up的后果。
• 不要只看平均值——平均值会被异常值拉偏。中位数更稳健。如果你有时间，可以同时看P50、P95、P99。
• 不要只测一个配置——grid/block大小对吞吐影响巨大。扫一遍参数空间，找到最佳点，再用这个点做后续对比。
• 不要同时改多个变量——控制变量法是科学实验的基础。一次只改一个，才能确定因果关系。

好了，这就是吞吐压测的方法论。说白了就是：warm-up、取中位数、扫参数、控制变量。把这四点做到位，你的测试结果就靠谱了。