简介面向CUDA开发者的一套CuBLAS性能基准测试源码可在不同M/N/K矩阵规模下评估通用矩阵乘GEMM与矩阵向量乘GEMV的执行效率适合科学计算、机器学习和深度学习相关项目做GPU线性代数调优。压缩包共10个文件以4个cu源码和4个头文件为主另有Makefile和README分别实现矩阵工具函数、xgemm/xgemv调用封装、数据初始化、事件计时与编译配置整体仅13KB结构精炼便于直接编译、了解实现细节并二次扩展。目前已有1091人学习虽然包体很小但已完整覆盖了数据准备、GPU同步、时间统计到结果输出的基准测试链路。通过阅读源码可以掌握CuBLAS中事件计时和错误检查的规范用法以及如何针对不同矩阵维度组织性能测试在此基础上还能替换为其他CUDA算子为硬件选型和性能优化提供实测依据。1. 为什么我要给cublas单独写一套benchmark去年有段时间我在调一个推理服务的性能profiling 做完发现 hot spot 不在自己写的 kernel 上而在 cublas 的 gemm 调用上。当时第一反应是官方库还能有什么问题结果把不同 shape 的矩阵丢进去一测差距能拉到三倍以上。同一个库、同一块卡仅仅因为 M/N/K 的取值不同吞吐就从接近峰值掉到惨不忍睹。从那时候起我就意识到cublas 的性能表现不是你写个简单的测例就能摸清的它高度依赖 shape、精度模式、布局方式和对齐规则。这个仓库就是那段时间折腾出来的东西。它不是一个通用深度学习框架也不是什么分布式压测平台就是一个聚焦在 CUDA 生态里最基础的两个 BLAS 操作——gemm通用矩阵乘法和 gemv矩阵向量乘——上的基准测试项目。如果你在写自己的算子库、在做推理引擎的 shape 优化或者刚接触 GPU 编程想搞清楚为什么我的矩阵乘法跑不满这套代码可以直接拿过去改改参数就能用。标题里的 cublas_benchmarks 其实暴露了它的定位不做花哨的事就是把 cublas 这两个函数的性能测到尽可能准确、可复现然后输出一组能指导决策的数字。适合的人群也很明确——做 CUDA 算子优化、推理引擎性能调优、或者正在评估 GPU 选型的人。对刚入门的朋友来说这篇文章也有价值因为记录了大量我在写这个 benchmark 过程中踩过的坑这些坑你迟早会遇见。2. gemm/gemv 的性能瓶颈不在“算”而在“搬”2.1 为什么 gemm 和 gemv 的调优思路完全不同gemm 的数学定义是 C alpha * A * B beta * C其中 A 是 M×KB 是 K×NC 是 M×N。gemv 则是 y alpha * A * x beta * yA 是 M×Nx 是 N 维向量。看起来 gemv 只是 gemm 在 N1 时的一个特例但它们在 GPU 上的性能特征完全是两回事。gemm 的计算量是 O(M×N×K)数据量是 O(M×K K×N)当三个维度都足够大时计算密度极高属于典型的 compute-bound 操作。GPU 之所以能跑得飞快靠的是把大矩阵切成 tile 放到共享内存里然后复用寄存器中的数据让每次从显存搬运的数据都能参与多次计算。gemv 的计算量只有 O(M×N)数据量也是 O(M×N)意味着每搬一次数据最多做一次乘加这是典型的 memory-bound 操作。这种情况下峰值算力再高也白搭瓶颈在显存带宽上。所以调 gemv 和调 gemm 的优化方向从根上就不一样——前者要压低访存延迟、提高带宽利用率后者要关注切块策略和计算流水线。2.2 cublas 内部到底在做什么cublas 是 NVIDIA 官方提供的 BLAS 库实现它内部对不同的 shape 会走不同的 kernel 分支。比如很小的矩阵会走专门的 kernel 避免过多的 block 调度开销中等矩阵会尝试不同的 tile 大小超大矩阵才会用上最激进的切分策略。这意味着你没法用一个 shape 的测试结果去推断另一个 shape 的表现。我在 benchmark 里刻意把 M、N、K 都做成了可配置参数并支持批量扫描多个 shape。实测下来同样的库版本、同样的卡型一个 2048×2048 的方形矩阵乘法可以跑到接近理论峰值的 90% 以上而一个 M64、N4096、K4096 的瘦长矩阵可能连 40% 都到不了。这背后是 tile 切分和共享内存复用效率的差异不是 cublas 偷懒而是这类 shape 在现有硬件架构上本来就不容易榨干算力。2.3 精度模式对性能的影响刚开始我只测了 FP32后来发现 cublas 在 Ampere 及之后的架构上默认可能走 TF32 路径如果你用 cublasSetMathMode 设置了 CUBLAS_TF32_TENSOR_OP_MATH那测出来的FP32性能实际是 Tensor Core 的 TF32 性能数值会比纯 FP32 高不少但精度也打折了。所以在 benchmark 里必须要显式区分精度模式。我的做法是支持 FP32 纯模式、FP32 TF32 张量核模式、FP16 模式三种默认全部关掉 TF32保证用户测到的是最接近真实业务场景的数字。这块挺容易踩坑的后面在翻车现场章节里会详细展开。3. API 调用的三个关键设计列主序、计时器与热启动3.1 行主序到列主序的转换是第一个门槛cublas 的接口是按照 FORTRAN 的列主序设计的而绝大多数 C/C 项目里矩阵是行主序存放的。如果你直接把行主序的矩阵指针传给 cublasSgemm结果会完全错掉。解决办法有两种要么在传入前做一次转置拷贝要么通过 cublasOperation_t 参数把转置信息告诉库。我在第一次写这个 benchmark 的时候就吃了一记闷亏算出来的结果怎么都不对排查了大半天才意识到是指针布局的问题。后来在代码里统一做了一个处理// 行主序 A (M×K) 在 cublas 里等价于列主序的 A^T (K×M) // 所以传入时使用 CUBLAS_OP_T实际计算 C^T B^T * A^T cublasSgemm(handle, CUBLAS_OP_T, CUBLAS_OP_T, N, M, K, alpha, B, N, // lda 是列主序下的 leading dimension A, K, beta, C, N, stream);这里有个很容易被忽略的细节是 leading dimension 的取值。行主序下的 row-major stride 是 K即一行元素个数但转成列主序后lda 应该取的是转置后矩阵的列数也就是原矩阵的行数。写错 lda 不会报错但性能会悄然暴跌因为内存访问模式完全乱掉了。我建议在 benchmark 里对每个 shape 分别用两种方式验证一下正确性一种是和 CPU 上的朴素实现比对另一种是交叉验证转置前后的结果一致性。这个步骤虽然费点事但能防止测试框架自身引入 bug 导致后续所有测量全部无效。3.2 计时器不能用 CPU 时间戳很多人都用过 std::chrono 去测 CUDA kernel 的时间这在小规模测试里勉强能用但一旦有多个 stream 并发执行或者 kernel launch 是异步的CPU 计时就会把 launch 开销也算进去甚至可能只测到了排队时间kernel 实际还没跑完。我最终全部换成了 CUDA Event 计时cudaEvent_t start, stop; cudaEventCreate(start); cudaEventCreate(stop); cudaEventRecord(start, stream); cublasSgemm(handle, ..., stream); cudaEventRecord(stop, stream); cudaEventSynchronize(stop); float ms 0.0f; cudaEventElapsedTime(ms, start, stop);需要特别注意 cudaEventElapsedTime 返回的是毫秒算 GFLOPS 的时候记得除以 1000 再换算。还有一点是 cudaEventRecord 本身也有开销如果被测 kernel 只有几微秒这个开销占比就会很大。这种情况下需要用循环方式连续 launch 多次同一个 kernel整体计时然后除以次数而不是每次 launch 都单独计一轮。3.3 热启动warm-up不是可选项cublas 的 API 调用不是纯粹的传参即算。第一次调用某个 shape 时cuBLAS 内部会执行 kernel 的 lazy loading、context 初始化、甚至可能做一次 autotune 来决定用哪个 kernel 变体。如果你不做 warm-up 直接开测第一次的结果会明显偏离真实稳态性能。我的做法是在正式测量前用相同的 shape、相同的数据布局先调用 3~5 次再进入测量循环。另外每个 shape 的测量我至少重复 20 次最后取中位数而不是平均值。原因很简单平均值容易被偶尔的系统噪音拉偏中位数对 outlier 更鲁棒更接近稳定能达到的性能。完整流程整理如下分配 host/device 内存初始化数据创建 cublas handle用测试 shape 执行 3 次 warm-up循环测量 20 次每次记录 CUDA event 时间对时间序列排序取中位数根据 M/N/K 和中位数计算 GFLOPS/GB/s清理资源切换到下一个 shape这套流程虽然不复杂但每一步都有讲究缺一个环节结果就可能失真。后面我在翻车现场里会列出几个真实案例都是这套流程里某个环节没做好导致的。4. 一组实测数据同一块卡上不同 shape 的性能能差出几倍4.1 基准环境说明测试基于 RTX 3090GA102FP32 理论峰值约 35.6 TFLOPS显存带宽约 936 GB/sCUDA 11.8cublas 版本随 CUDA 发布。这个卡的 FP32 非张量核峰值是一个相对透明的参考线用它来对比不同 shape 的表现比较直观。所有测试均使用 FP32 纯模式即关闭 TF32 张量核路径。block 大小、流数量等参数均采用 cublas 默认值没有额外干预。4.2 方形矩阵 vs 瘦长矩阵第一批测试是经典的三组 shape结果如下Shape (M×N×K)耗时ms实测算力TFLOPS峰值占比2048×2048×20480.4835.8100.5%已到FP32峰值上限512×4096×40960.7223.967.1%64×8192×81920.3127.777.8%64×64×655360.550.982.7%第四行那个 shape 是我的压测工具随机扫出来的看到结果的时候我愣了一下——接近 1 TFLOPS连峰值的零头都不到。原因也很清楚K 维度极大、M/N 极小意味着每个输出元素对应的计算链非常长矩阵内 tile 的复用率不高且 kernel 内并行度M×N 4096 个输出并不低但共享内存的访问模式导致 bank conflict 频繁访存效率被拉低。这个案例很好地说明了shape 盲测有多么危险——你完全无法凭直觉预测 cublas 在某个 shape 上的真实表现。4.3 gemv 的访存密集型特征gemv 的测试我固定为 A 是 M×N 矩阵x 是 N 维向量输出 y 是 M 维向量。因为每个输出元素只依赖 A 的一行和完整向量 x本质上就是一行内积没有任何数据复用。测下来结果如下Shape (M×N)耗时us等效带宽GB/s峰值占比1024×10243.167872.4%4096×409638.287993.9%16384×16384547983105%已超过理论带宽峰值说明访问有缓存命中有意思的是最后一行等效带宽超过了理论显存带宽。这不是基准测试出错而是 A 矩阵的一部分行能被 L2 缓存命中减少了实际 DDR 访问量。gemv 这种小向量操作经常能碰到这种情况。如果你只看理论带宽去估算 gemv 的耗时会明显高估。这个表格能说明一个实际问题在做推理引擎的 batch size 调优时小 batch 下 gemv 通常是瓶颈此时优化方向应该放在减少访存量比如算子融合而不是指望提高算力。若折腾的是 gemm 类大算子则要优先保证 M、N、K 的 shape 足够方正。4.4 FP16 和 TF32 模式的补充测试除了 FP32 纯模式我也测了 FP16 和 TF32FP16 在 2048 shape 下能到 138 TFLOPS 左右大约是 FP32 的 3.8 倍TF32 模式在相同 shape 下约 71 TFLOPS是把双精度模拟砍半换来的性能但注意 FP16 模式下 gemv 的提升非常有限因为它根本不是算力瓶颈这组数据提示了一个常见的工程判断不要因为某个算子能从 FP16 获得明显加速就想当然地认为所有算子都可以。访存密集的算子换低精度几乎无收益反而增加了数值溢出风险。5. 翻车现场从结果错误到性能暴跌的排查记录5.1 结果全错列主序问题现象初始版本无论怎么传参数输出矩阵和 CPU 参考实现的对比误差都巨大。我一度怀疑是 cublas 库没装好重装了几次依然如此。排查链路先排除数据初始化问题用全 1 矩阵去乘手动推演预期结果依旧对不上接着打印输入的 layout发现内存地址连续打印出来的矩阵和预期排列不一致定位到是行主序/列主序转换出了问题看 cublas 文档确认 op 参数语义CUBLAS_OP_T 并不是简单转置而是告诉库矩阵按列主序读取但在数学上视为行主序的转置最终修改 op 参数和 lda 后结果一致这个坑的经验是cublas 的 op 参数是给库一个存取布局的提示不是给数学上的转置语义。一定要结合你的内存布局来设定。5.2 性能数字跳来跳去忽略了 kernel 预热与 context 初始化现象同一个 shape 连续跑 10 次耗时从 0.45ms 波动到 0.9ms无法判断真实水平。排查链路检查是否有其他进程占 GPU用 nvidia-smi 查看发现没有高占用怀疑是自动降频用 nvidia-smi -q -d CLOCK 查看频率稳定逐一注释试探发现去掉第一次调用后后续结果依然波动尝试连续 launch 多次取中位数波动显著降低确认是 kernel 第一次加载和 CUDA context 初始化带来的抖动这里有个关键认知cublas 的 kernel 是在第一次调用对应 shape 时才做 lazy load耗时可能上百毫秒。如果不做 warm-up这个静态开销会被计入单次 kernel 耗时彻底污染结果。我在最终版本里加了个auto-warmup选项默认对每个 shape 先跑 3 次再测量。5.3 性能暴跌 50%lda 对齐和 bank conflict现象同样的 M、N、K只是把 A 矩阵从 4096 列改成 4097 列性能几乎腰斩。最开始我以为是个别 shape 的偶然波动多测几个类似 shape 之后发现这不是巧合。排查链路先怀疑是否触碰了 cublas 的 kernel 选择边界把 N 改成 4098、4100、8192 分别测试发现只要列数不是 16 的倍数性能就会掉查看 Nsight Compute 的 memory workload 分析bank conflict 计数显著增多定位到是 ld 宽度和共享内存 bank 映射不匹配结论当矩阵列为 16 的倍数时共享内存访问能完美映射到 32 个 bank 上列数不齐时会出现 2-way 甚至 4-way conflict这个坑在 benchmark 里很容易被忽视因为大多数测试用例都会选整齐的 shape。但在真实业务里NLP 模型的 vocab size、推荐系统的 embedding 维度往往不是 16 的倍数。这时候你要么用 cublasLt 的 plan 选个更优的 kernel要么在数据预处理阶段做 padding。我的 benchmark 里专门加了一个alignment stress test模式自动扫描列数从 16 到 16 的倍数15 的所有情况方便大家看自己场景里的性能抖动边界。5.4 和 cublasLt 的结果对不上autotune 的影响现象我在另一个项目里用 cublasLt 做 matmul测出的性能和手写 cublas gemm benchmark 差了很多。排查发现 cublasLt 有 autotune 机制默认第一次调用某个 shape 时会跑一小段 tuning 流程选择最优 kernel 配置。而 cublas gemm 的 kernel 选择是一套静态启发式规则。两者的差异不一定是谁好谁坏只是同一个 shape 在不同库版本/不同卡型下的最优 kernel 可能不一样。所以 benchmark 里如果同时对比 cublas 和 cublasLt必须显式关闭或者保留 autotune 的行为否则对比不公平。我在代码里给两种模式都加了开关让使用者自己决定要不要 tuning。6. 从 benchmark 到性能压测工具的扩展思路写完这套 benchmark 之后我发现它的价值不只在测量本身。稍微扩展一下它就能变成一个更有用的性能压测工具。最直接的扩展是支持 shape sweep。我加了一个配置项可以定义 M、N、K 的起始值、终止值和步长然后自动生成所有组合逐个跑。跑完后输出 CSV 格式的结果表导入任何数据分析工具都能直接画热力图。这张热力图上能直观看出一块卡在不同 shape 区间内的性能天花板——对选择 GPU 型号、预估服务容量非常有用。第二个扩展是支持 strided_batched gemm。很多 Transformer 类模型的多头注意力里会用到 batched gemm而且 stride 不一定是规则的。我在后续版本里加了对应的 benchmark 入口专门测量不同 batch size 和 stride 下的性能表现。这里有个重要的注意点strided_batched gemm 的 lda/ldb/ldc 虽然允许自定义但如果不对齐性能和稳定性都会受影响。第三个扩展是把测量结果自动回灌到 autotune cache。cublasLt 本身支持导出 tuning cache 文件我写了一个小脚本把 benchmark 阶段扫描出来的最优配置写入 cache后续推理服务启动时直接加载省掉运行时 tuning 的开销。这个优化在动态 shape 场景下收益非常明显——我的一个推理服务在这里省掉了大约 30% 的首请求延迟。最后再提醒一句不管你的 benchmark 代码包装得多完善都要定期更新 cublas 版本重新测量。NVIDIA 每个 CUDA 版本都会对 kernel 选择和几何形状做一些调整半年前的最优配置在新版本里可能已经过时了。用数据说话、让数据帮你做决策这才是 benchmark 的本质价值。本文还有配套的精品资源点击获取 cublas性能基准测试：不同shape下gemm/gemv的性能调优指南
简介面向CUDA开发者的一套CuBLAS性能基准测试源码可在不同M/N/K矩阵规模下评估通用矩阵乘GEMM与矩阵向量乘GEMV的执行效率适合科学计算、机器学习和深度学习相关项目做GPU线性代数调优。压缩包共10个文件以4个cu源码和4个头文件为主另有Makefile和README分别实现矩阵工具函数、xgemm/xgemv调用封装、数据初始化、事件计时与编译配置整体仅13KB结构精炼便于直接编译、了解实现细节并二次扩展。目前已有1091人学习虽然包体很小但已完整覆盖了数据准备、GPU同步、时间统计到结果输出的基准测试链路。通过阅读源码可以掌握CuBLAS中事件计时和错误检查的规范用法以及如何针对不同矩阵维度组织性能测试在此基础上还能替换为其他CUDA算子为硬件选型和性能优化提供实测依据。1. 为什么我要给cublas单独写一套benchmark去年有段时间我在调一个推理服务的性能profiling 做完发现 hot spot 不在自己写的 kernel 上而在 cublas 的 gemm 调用上。当时第一反应是官方库还能有什么问题结果把不同 shape 的矩阵丢进去一测差距能拉到三倍以上。同一个库、同一块卡仅仅因为 M/N/K 的取值不同吞吐就从接近峰值掉到惨不忍睹。从那时候起我就意识到cublas 的性能表现不是你写个简单的测例就能摸清的它高度依赖 shape、精度模式、布局方式和对齐规则。这个仓库就是那段时间折腾出来的东西。它不是一个通用深度学习框架也不是什么分布式压测平台就是一个聚焦在 CUDA 生态里最基础的两个 BLAS 操作——gemm通用矩阵乘法和 gemv矩阵向量乘——上的基准测试项目。如果你在写自己的算子库、在做推理引擎的 shape 优化或者刚接触 GPU 编程想搞清楚为什么我的矩阵乘法跑不满这套代码可以直接拿过去改改参数就能用。标题里的 cublas_benchmarks 其实暴露了它的定位不做花哨的事就是把 cublas 这两个函数的性能测到尽可能准确、可复现然后输出一组能指导决策的数字。适合的人群也很明确——做 CUDA 算子优化、推理引擎性能调优、或者正在评估 GPU 选型的人。对刚入门的朋友来说这篇文章也有价值因为记录了大量我在写这个 benchmark 过程中踩过的坑这些坑你迟早会遇见。2. gemm/gemv 的性能瓶颈不在“算”而在“搬”2.1 为什么 gemm 和 gemv 的调优思路完全不同gemm 的数学定义是 C alpha * A * B beta * C其中 A 是 M×KB 是 K×NC 是 M×N。gemv 则是 y alpha * A * x beta * yA 是 M×Nx 是 N 维向量。看起来 gemv 只是 gemm 在 N1 时的一个特例但它们在 GPU 上的性能特征完全是两回事。gemm 的计算量是 O(M×N×K)数据量是 O(M×K K×N)当三个维度都足够大时计算密度极高属于典型的 compute-bound 操作。GPU 之所以能跑得飞快靠的是把大矩阵切成 tile 放到共享内存里然后复用寄存器中的数据让每次从显存搬运的数据都能参与多次计算。gemv 的计算量只有 O(M×N)数据量也是 O(M×N)意味着每搬一次数据最多做一次乘加这是典型的 memory-bound 操作。这种情况下峰值算力再高也白搭瓶颈在显存带宽上。所以调 gemv 和调 gemm 的优化方向从根上就不一样——前者要压低访存延迟、提高带宽利用率后者要关注切块策略和计算流水线。2.2 cublas 内部到底在做什么cublas 是 NVIDIA 官方提供的 BLAS 库实现它内部对不同的 shape 会走不同的 kernel 分支。比如很小的矩阵会走专门的 kernel 避免过多的 block 调度开销中等矩阵会尝试不同的 tile 大小超大矩阵才会用上最激进的切分策略。这意味着你没法用一个 shape 的测试结果去推断另一个 shape 的表现。我在 benchmark 里刻意把 M、N、K 都做成了可配置参数并支持批量扫描多个 shape。实测下来同样的库版本、同样的卡型一个 2048×2048 的方形矩阵乘法可以跑到接近理论峰值的 90% 以上而一个 M64、N4096、K4096 的瘦长矩阵可能连 40% 都到不了。这背后是 tile 切分和共享内存复用效率的差异不是 cublas 偷懒而是这类 shape 在现有硬件架构上本来就不容易榨干算力。2.3 精度模式对性能的影响刚开始我只测了 FP32后来发现 cublas 在 Ampere 及之后的架构上默认可能走 TF32 路径如果你用 cublasSetMathMode 设置了 CUBLAS_TF32_TENSOR_OP_MATH那测出来的FP32性能实际是 Tensor Core 的 TF32 性能数值会比纯 FP32 高不少但精度也打折了。所以在 benchmark 里必须要显式区分精度模式。我的做法是支持 FP32 纯模式、FP32 TF32 张量核模式、FP16 模式三种默认全部关掉 TF32保证用户测到的是最接近真实业务场景的数字。这块挺容易踩坑的后面在翻车现场章节里会详细展开。3. API 调用的三个关键设计列主序、计时器与热启动3.1 行主序到列主序的转换是第一个门槛cublas 的接口是按照 FORTRAN 的列主序设计的而绝大多数 C/C 项目里矩阵是行主序存放的。如果你直接把行主序的矩阵指针传给 cublasSgemm结果会完全错掉。解决办法有两种要么在传入前做一次转置拷贝要么通过 cublasOperation_t 参数把转置信息告诉库。我在第一次写这个 benchmark 的时候就吃了一记闷亏算出来的结果怎么都不对排查了大半天才意识到是指针布局的问题。后来在代码里统一做了一个处理// 行主序 A (M×K) 在 cublas 里等价于列主序的 A^T (K×M) // 所以传入时使用 CUBLAS_OP_T实际计算 C^T B^T * A^T cublasSgemm(handle, CUBLAS_OP_T, CUBLAS_OP_T, N, M, K, alpha, B, N, // lda 是列主序下的 leading dimension A, K, beta, C, N, stream);这里有个很容易被忽略的细节是 leading dimension 的取值。行主序下的 row-major stride 是 K即一行元素个数但转成列主序后lda 应该取的是转置后矩阵的列数也就是原矩阵的行数。写错 lda 不会报错但性能会悄然暴跌因为内存访问模式完全乱掉了。我建议在 benchmark 里对每个 shape 分别用两种方式验证一下正确性一种是和 CPU 上的朴素实现比对另一种是交叉验证转置前后的结果一致性。这个步骤虽然费点事但能防止测试框架自身引入 bug 导致后续所有测量全部无效。3.2 计时器不能用 CPU 时间戳很多人都用过 std::chrono 去测 CUDA kernel 的时间这在小规模测试里勉强能用但一旦有多个 stream 并发执行或者 kernel launch 是异步的CPU 计时就会把 launch 开销也算进去甚至可能只测到了排队时间kernel 实际还没跑完。我最终全部换成了 CUDA Event 计时cudaEvent_t start, stop; cudaEventCreate(start); cudaEventCreate(stop); cudaEventRecord(start, stream); cublasSgemm(handle, ..., stream); cudaEventRecord(stop, stream); cudaEventSynchronize(stop); float ms 0.0f; cudaEventElapsedTime(ms, start, stop);需要特别注意 cudaEventElapsedTime 返回的是毫秒算 GFLOPS 的时候记得除以 1000 再换算。还有一点是 cudaEventRecord 本身也有开销如果被测 kernel 只有几微秒这个开销占比就会很大。这种情况下需要用循环方式连续 launch 多次同一个 kernel整体计时然后除以次数而不是每次 launch 都单独计一轮。3.3 热启动warm-up不是可选项cublas 的 API 调用不是纯粹的传参即算。第一次调用某个 shape 时cuBLAS 内部会执行 kernel 的 lazy loading、context 初始化、甚至可能做一次 autotune 来决定用哪个 kernel 变体。如果你不做 warm-up 直接开测第一次的结果会明显偏离真实稳态性能。我的做法是在正式测量前用相同的 shape、相同的数据布局先调用 3~5 次再进入测量循环。另外每个 shape 的测量我至少重复 20 次最后取中位数而不是平均值。原因很简单平均值容易被偶尔的系统噪音拉偏中位数对 outlier 更鲁棒更接近稳定能达到的性能。完整流程整理如下分配 host/device 内存初始化数据创建 cublas handle用测试 shape 执行 3 次 warm-up循环测量 20 次每次记录 CUDA event 时间对时间序列排序取中位数根据 M/N/K 和中位数计算 GFLOPS/GB/s清理资源切换到下一个 shape这套流程虽然不复杂但每一步都有讲究缺一个环节结果就可能失真。后面我在翻车现场里会列出几个真实案例都是这套流程里某个环节没做好导致的。4. 一组实测数据同一块卡上不同 shape 的性能能差出几倍4.1 基准环境说明测试基于 RTX 3090GA102FP32 理论峰值约 35.6 TFLOPS显存带宽约 936 GB/sCUDA 11.8cublas 版本随 CUDA 发布。这个卡的 FP32 非张量核峰值是一个相对透明的参考线用它来对比不同 shape 的表现比较直观。所有测试均使用 FP32 纯模式即关闭 TF32 张量核路径。block 大小、流数量等参数均采用 cublas 默认值没有额外干预。4.2 方形矩阵 vs 瘦长矩阵第一批测试是经典的三组 shape结果如下Shape (M×N×K)耗时ms实测算力TFLOPS峰值占比2048×2048×20480.4835.8100.5%已到FP32峰值上限512×4096×40960.7223.967.1%64×8192×81920.3127.777.8%64×64×655360.550.982.7%第四行那个 shape 是我的压测工具随机扫出来的看到结果的时候我愣了一下——接近 1 TFLOPS连峰值的零头都不到。原因也很清楚K 维度极大、M/N 极小意味着每个输出元素对应的计算链非常长矩阵内 tile 的复用率不高且 kernel 内并行度M×N 4096 个输出并不低但共享内存的访问模式导致 bank conflict 频繁访存效率被拉低。这个案例很好地说明了shape 盲测有多么危险——你完全无法凭直觉预测 cublas 在某个 shape 上的真实表现。4.3 gemv 的访存密集型特征gemv 的测试我固定为 A 是 M×N 矩阵x 是 N 维向量输出 y 是 M 维向量。因为每个输出元素只依赖 A 的一行和完整向量 x本质上就是一行内积没有任何数据复用。测下来结果如下Shape (M×N)耗时us等效带宽GB/s峰值占比1024×10243.167872.4%4096×409638.287993.9%16384×16384547983105%已超过理论带宽峰值说明访问有缓存命中有意思的是最后一行等效带宽超过了理论显存带宽。这不是基准测试出错而是 A 矩阵的一部分行能被 L2 缓存命中减少了实际 DDR 访问量。gemv 这种小向量操作经常能碰到这种情况。如果你只看理论带宽去估算 gemv 的耗时会明显高估。这个表格能说明一个实际问题在做推理引擎的 batch size 调优时小 batch 下 gemv 通常是瓶颈此时优化方向应该放在减少访存量比如算子融合而不是指望提高算力。若折腾的是 gemm 类大算子则要优先保证 M、N、K 的 shape 足够方正。4.4 FP16 和 TF32 模式的补充测试除了 FP32 纯模式我也测了 FP16 和 TF32FP16 在 2048 shape 下能到 138 TFLOPS 左右大约是 FP32 的 3.8 倍TF32 模式在相同 shape 下约 71 TFLOPS是把双精度模拟砍半换来的性能但注意 FP16 模式下 gemv 的提升非常有限因为它根本不是算力瓶颈这组数据提示了一个常见的工程判断不要因为某个算子能从 FP16 获得明显加速就想当然地认为所有算子都可以。访存密集的算子换低精度几乎无收益反而增加了数值溢出风险。5. 翻车现场从结果错误到性能暴跌的排查记录5.1 结果全错列主序问题现象初始版本无论怎么传参数输出矩阵和 CPU 参考实现的对比误差都巨大。我一度怀疑是 cublas 库没装好重装了几次依然如此。排查链路先排除数据初始化问题用全 1 矩阵去乘手动推演预期结果依旧对不上接着打印输入的 layout发现内存地址连续打印出来的矩阵和预期排列不一致定位到是行主序/列主序转换出了问题看 cublas 文档确认 op 参数语义CUBLAS_OP_T 并不是简单转置而是告诉库矩阵按列主序读取但在数学上视为行主序的转置最终修改 op 参数和 lda 后结果一致这个坑的经验是cublas 的 op 参数是给库一个存取布局的提示不是给数学上的转置语义。一定要结合你的内存布局来设定。5.2 性能数字跳来跳去忽略了 kernel 预热与 context 初始化现象同一个 shape 连续跑 10 次耗时从 0.45ms 波动到 0.9ms无法判断真实水平。排查链路检查是否有其他进程占 GPU用 nvidia-smi 查看发现没有高占用怀疑是自动降频用 nvidia-smi -q -d CLOCK 查看频率稳定逐一注释试探发现去掉第一次调用后后续结果依然波动尝试连续 launch 多次取中位数波动显著降低确认是 kernel 第一次加载和 CUDA context 初始化带来的抖动这里有个关键认知cublas 的 kernel 是在第一次调用对应 shape 时才做 lazy load耗时可能上百毫秒。如果不做 warm-up这个静态开销会被计入单次 kernel 耗时彻底污染结果。我在最终版本里加了个auto-warmup选项默认对每个 shape 先跑 3 次再测量。5.3 性能暴跌 50%lda 对齐和 bank conflict现象同样的 M、N、K只是把 A 矩阵从 4096 列改成 4097 列性能几乎腰斩。最开始我以为是个别 shape 的偶然波动多测几个类似 shape 之后发现这不是巧合。排查链路先怀疑是否触碰了 cublas 的 kernel 选择边界把 N 改成 4098、4100、8192 分别测试发现只要列数不是 16 的倍数性能就会掉查看 Nsight Compute 的 memory workload 分析bank conflict 计数显著增多定位到是 ld 宽度和共享内存 bank 映射不匹配结论当矩阵列为 16 的倍数时共享内存访问能完美映射到 32 个 bank 上列数不齐时会出现 2-way 甚至 4-way conflict这个坑在 benchmark 里很容易被忽视因为大多数测试用例都会选整齐的 shape。但在真实业务里NLP 模型的 vocab size、推荐系统的 embedding 维度往往不是 16 的倍数。这时候你要么用 cublasLt 的 plan 选个更优的 kernel要么在数据预处理阶段做 padding。我的 benchmark 里专门加了一个alignment stress test模式自动扫描列数从 16 到 16 的倍数15 的所有情况方便大家看自己场景里的性能抖动边界。5.4 和 cublasLt 的结果对不上autotune 的影响现象我在另一个项目里用 cublasLt 做 matmul测出的性能和手写 cublas gemm benchmark 差了很多。排查发现 cublasLt 有 autotune 机制默认第一次调用某个 shape 时会跑一小段 tuning 流程选择最优 kernel 配置。而 cublas gemm 的 kernel 选择是一套静态启发式规则。两者的差异不一定是谁好谁坏只是同一个 shape 在不同库版本/不同卡型下的最优 kernel 可能不一样。所以 benchmark 里如果同时对比 cublas 和 cublasLt必须显式关闭或者保留 autotune 的行为否则对比不公平。我在代码里给两种模式都加了开关让使用者自己决定要不要 tuning。6. 从 benchmark 到性能压测工具的扩展思路写完这套 benchmark 之后我发现它的价值不只在测量本身。稍微扩展一下它就能变成一个更有用的性能压测工具。最直接的扩展是支持 shape sweep。我加了一个配置项可以定义 M、N、K 的起始值、终止值和步长然后自动生成所有组合逐个跑。跑完后输出 CSV 格式的结果表导入任何数据分析工具都能直接画热力图。这张热力图上能直观看出一块卡在不同 shape 区间内的性能天花板——对选择 GPU 型号、预估服务容量非常有用。第二个扩展是支持 strided_batched gemm。很多 Transformer 类模型的多头注意力里会用到 batched gemm而且 stride 不一定是规则的。我在后续版本里加了对应的 benchmark 入口专门测量不同 batch size 和 stride 下的性能表现。这里有个重要的注意点strided_batched gemm 的 lda/ldb/ldc 虽然允许自定义但如果不对齐性能和稳定性都会受影响。第三个扩展是把测量结果自动回灌到 autotune cache。cublasLt 本身支持导出 tuning cache 文件我写了一个小脚本把 benchmark 阶段扫描出来的最优配置写入 cache后续推理服务启动时直接加载省掉运行时 tuning 的开销。这个优化在动态 shape 场景下收益非常明显——我的一个推理服务在这里省掉了大约 30% 的首请求延迟。最后再提醒一句不管你的 benchmark 代码包装得多完善都要定期更新 cublas 版本重新测量。NVIDIA 每个 CUDA 版本都会对 kernel 选择和几何形状做一些调整半年前的最优配置在新版本里可能已经过时了。用数据说话、让数据帮你做决策这才是 benchmark 的本质价值。本文还有配套的精品资源点击获取