简介cublas_benchmarks是一个用于测试通用矩阵乘法GEMM和矩阵向量乘法GEMV性能的基准项目主要面向CUDA开发者。项目通过调整矩阵和向量的规模参数调用cuBLAS库中的对应函数精确记录每次运算的耗时帮助使用者量化图形处理器在不同计算任务下的表现也可用于对比不同硬件或驱动环境带来的性能差异为深度学习、科学计算等场景的算法选型提供依据。在训练神经网络时矩阵乘法是反向传播中的核心操作而矩阵向量乘法常见于线性回归等场景因此该基准工具能帮助开发者发现性能瓶颈。压缩包共10个文件包含4个头文件、4个CUDA源文件、1个Makefile和1个README整体仅13KB头文件负责数据结构和函数声明源文件实现数据初始化、性能测试和计时逻辑Makefile可辅助编译README则说明代码结构。该资源已有1095人学习下载适合希望优化矩阵运算效率、深入理解cuBLAS调用方式的C开发者参考也可作为搭建性能测试框架的起点。1. cublas_benchmarks 到底在测什么做推理服务优化时我经常遇到这种情况某个算子耗时居高不下分析工具指向矩阵乘法但拿不出一个可信的数字来证明它该快还是该慢。cublas_benchmarks 就是干这个的——它把 cuBLAS 里的 GEMM通用矩阵乘法和 GEMV矩阵向量乘单独拎出来用可控的 M、N、K 维度组合跑一遍把每次调用的时间记录成 TFLOPS 和带宽数据。适合三类人做算子性能回归的、对比不同 GPU 或驱动版本的、以及刚接触 CUDA 想搞清楚自己显卡算力上限的开发者。反直觉的一点是很多新手跑完第一轮数据都会怀疑人生小矩阵的 GEMV 只能跑出个位数 TFLOPS这不是库的问题是访存瓶颈在作祟。2. 为什么 GEMM 和 GEMV 值得单独做基准2.1 GEMM 与 GEMV 的计算画像差异GEMM 的经典定义是 C alpha * A * B beta * C其中 A 是 M×KB 是 K×NC 是 M×N。把 C 矩阵算出来需要约 2 * M * N * K 次浮点运算而数据量只有 M×K K×N M×N 个元素。在 M、N、K 都大于 1024 时计算量和访存量不在一个量级所以 GEMM 是典型的计算密集型操作理论上限是 GPU 的峰值浮点算力。GEMV 则是 y alpha * A * x beta * yA 是 M×Nx 是 N 维向量。总计算量只有 2 * M * N 次却要把 A 矩阵的全部 M×N 个元素从显存搬出来。每读一个元素只做两次运算这时候算力完全闲置瓶颈在显存带宽上。所以跑 GEMV 时盯着 TFLOPS 没有意义应该看等效带宽GB/s有没有接近硬件上限。这个差别决定了 benchmark 的设计同一个工具里GEMM 和 GEMV 要输出不同的主指标混在一起分析会得出错误结论。我看到过有人用 GEMV 的低 TFLOPS 数值去质疑 cuBLAS 的性能其实是没分清计算密集和访存密集。2.2 用 cuBLAS 而不是自己写 kernel 的理由自己做矩阵乘法 kernel 并不难难的是在各种 shape 下都逼近硬件上限。cuBLAS 在每个 GPU 架构上都做了 kernel 选择cudaCore 的普通 kernel、Tensor Core 的专用 kernel还会根据 M、N、K 的整除性、对齐方式、batch 大小来决定用哪个切分策略。自己写一个 kernel 跑某个固定维度可能不输 cuBLAS但换一个非方形 shape 立刻现原形。benchmark 的目的是量化「当前环境下的真实上限」不是展示手写 kernel 的极限。所以用 cuBLAS 作为被测对象是合理的它是大多数上层框架实际调用的库测它就是在测生产环境的真实路径。另一个实际原因是 cuBLAS 的处理逻辑本身就是黑匣子基准测试是观察它的手段不需要去研究它的实现细节。2.3 benchmark 的核心指标定义GEMM 的核心指标是 TFLOPS计算公式是 2 * M * N * K / 时间/s / 1e12。这里的 2 来自乘法和加法各一次。GEMV 的核心指标是等效带宽需要读取的字节数约为 M * N * 4float 类型再算上写回 y 向量的 M * 4 字节总共大约是M * N M* 4 字节用 字节数 / 时间 / 1e9 得到 GB/s。另外有一个容易忽略的指标与理论峰值的百分比。拿 GEMM 的实测 TFLOPS 除以显卡的理论 TFLOPSfloat 精度下得到一个效率百分比。业界经验是方形大矩阵能跑到 85% 以上就算正常明显低于这个值就要检查是不是没关 ECC、驱动降频或散热问题。理论峰值可以从设备属性里拿到 SM 数量和时钟频率大致估算不必精读汇编指令。3. 把 cublas_benchmarks 跑起来依赖、编译与最小命令行3.1 环境检查与编译流程常见的做法是先确认两件事CUDA 工具链版本和显卡算力。cuBLAS 是随 CUDA toolkit 一起发布的动态库如果你的程序用了 cuBLAS运行时还需要确保动态库路径正确。执行下面的命令做前置检查nvcc --version nvidia-smi --query-gpuname,compute_cap,memory.total --formatcsv第一行看编译器版本第二行看驱动识别的显卡名、算力代数和显存总量。算力代数是 8.0 以上意味着有 Tensor Core跑 benchmark 时可以留意 TF32 路径是否生效。某些系统里 nvcc 版本和驱动版本不一致如果编译报找不到 cublas 头文件优先检查 CUDA 路径是否加入了环境变量。我一般会在 ~/.bashrc 里显式 export 这些变量避免系统里同时装多套 CUDA 时张冠李戴。CMake 是最省事的构建方式。一个最小 CMakeLists.txt 大致长这样cmake_minimum_required(VERSION 3.18) project(cublas_benchmarks LANGUAGES CXX CUDA) find_package(CUDAToolkit REQUIRED) add_executable(cublas_benchmarks main.cpp) target_link_libraries(cublas_benchmarks PRIVATE CUDA::cublas CUDA::cudart) set_target_properties(cublas_benchmarks PROPERTIES CUDA_ARCHITECTURES 80 )这里的关键是find_package(CUDAToolkit REQUIRED)它会自动定位 CUDA 头文件和库文件。CUDA_ARCHITECTURES指定了编译目标架构80 对应 Ampere 架构如果你在更老的卡上跑需要改成对应的算力数字否则运行时会因为 SASS 不兼容报错。还有一点容易被忽略编译选项里加上-O2不要开-G调试模式调试模式下 kernel 性能会差一个量级测出来的数据完全没有参考价值。编译命令也很简单mkdir -p build cd build cmake .. -DCMAKE_BUILD_TYPERelease make -j$(nproc)编译完成后得到一个可执行文件。如果链接时报找不到 cublas说明 CMake 没有正确找到 CUDA 库路径可以手动指定-DCUDAToolkit_ROOT/path/to/cuda再试。3.2 最小命令行与输出字段这个 benchmark 的最小用法是传入一组维度参数。参数设计参考了常见的矩阵乘法测试工具的约定并不复杂./cublas_benchmarks --m 4096 --n 4096 --k 4096 --iter 100 --warmup 10上面这条命令会运行一个 4096×4096×4096 的 GEMM先做 10 次预热再正式跑 100 次取平均值最后把结果打印出来。输出大致是这种格式shape: M4096 N4096 K4096 alpha1.00 beta0.00 avg_time: 0.623 ms min_time: 0.601 ms TFLOPS: 22.06 efficiency: 91.2%为什么同时打印 avg_time 和 min_time 很重要因为 GPU 执行期间可能被其他任务打断平均值会被偶发噪声抬高最小值更接近真实算力上限。我通常建议看最小值它反映的是 GPU 全力工作时的极限而平均值是你实际跑负载时可能遇到的表现。3.3 参数怎么设维度、循环次数、warmup 与数据类型下表整理了关键参数的含义和建议初始值这个表已经够覆盖绝大多数使用场景。参数含义建议初始值说明--mA 矩阵行数4096影响 kernel 切分不同量级差异极大--nB 矩阵列数4096与 m 组合扫描时优先固定一个--k内部分解维度4096越大越接近计算密集--iter正式计时循环次数100少于 30 次误差较大--warmup预热次数10少于 5 次结果无意义--dtype精度float可选 double / half--tensor-core是否启用 Tensor Core 路径auto对 Ampere 及以上架构有效--gemv切换到 GEMV 模式无测矩阵向量乘时使用warmup 次数为什么是 10 次起步cuBLAS 第一次调用会触发一些 lazy 初始化比如加载 heuristics 表、cudaFree 缓存分配、kernel 模块加载这些开销可能达到几百微秒直接淹没掉小矩阵本身的执行时间。10 次预热后这些一次性开销基本消失后续数据才稳定。iter 次数设太低时钟粒度导致的噪音占比会变大设太高大矩阵跑一次都几毫秒100 次也就是几百毫秒完全无压力所以建议往大设。GEMV 模式一般这么跑./cublas_benchmarks --gemv --m 10000 --n 4096 --iter 1000 --warmup 50因为 GEMV 单次执行极快即使 1000 次循环也只是一眨眼的事iter 可以放心放大。GEMV 是访存密集不具备 GEMM 的复用的优势预热次数多反而能让数据更平滑。4. 读懂输出用实际数据对齐理论峰值4.1 从时间到算力的换算逻辑跑完 benchmark 手里只有一个耗时要让它有意义就需要换算。GEMM 的 TFLOPS 计算我在前面已经给过公式这里补充它为什么是 2 乘而不是 3 乘或更复杂。矩阵乘法里C 矩阵的每个元素是 A 的一行和 B 的一列做点积长度是 K每个点积需要 K 次乘法和 K 次加法所以每个 C 元素对应 2K 次操作。整个 C 有 MN 个元素总操作数就是 2MNK。alpha 和 beta 的缩放加减属于额外开销在 K 很大的情况下占比可忽略基准工具通常不把它计入峰值。带宽换算就是另一套逻辑了。GEMV 模式读取的 A 矩阵大小是 MNsizeof(float)读一次 x 向量N4 字节写一次 y 向量M4 字节。忽略 x 向量的读取等效带宽约等于 MN4 / 时间。有些实现会报告「有效带宽」只算 A 矩阵的读取有的会报告「总带宽」把这些都加上不同的工具含义不同比较前先确认口径。4.2 同一块卡在不同 shape 下的巨大差异最典型的一个发现是同样是 1000 万次浮点操作方形大矩阵和瘦长矩阵的 TFLOPS 能差 10 倍以上。拿模拟数据举例在一张中等定位的显卡上MNK4096 可能跑到 20 TFLOPS而 M1024N64K4096 可能只有 4 TFLOPS 左右。原因有两个一是瘦长矩阵的切分被迫使用更小的 tile无法把 Tensor Core 的 16×16×16 计算块打满二是很小的 N 会让 C 矩阵的写回访存次数按 tile 放大访存占比上升。GEMV 的差异更多来自矩阵的行数和列数比例。A 矩阵按行主序存储时行数大意味着 y 向量的写回次数多写回可能遇到非连续地址。某个测试里 M1,000,000N16 的 GEMV 跑到的带宽不到 M100,000N256 的一半就是因为 y 向量的写回规模和缓存命中率都变了。4.3 一个可信的基准输出应该长什么样为了让结果能进报告或做回归对比工具的输出别只打印正文里的那种格式化文本还应该输出 CSV 行。这样同类数据可以直接叠加到同一个表格里。推荐的字段有时间戳、GPU 名称、CUDA 版本、op_typeGEMM 或 GEMV、M、N、K、dtype、Tensor Core 开关状态、耗时毫秒、TFLOPS 和等效带宽。op_type,m,n,k,dtype,tensor_core,avg_ms,min_ms,tflops gemm,4096,4096,4096,float,on,0.623,0.601,22.06 gemm,4096,1024,4096,float,on,0.291,0.287,11.85 gemv,10000,4096,float,on,0.018,0.017,4.52注意 GEMV 行的维度字段只有两个因为不存在 K 这个维度。输出 CSV 时把 min_ms 和 avg_ms 都带上既能看到稳定性也能看到是否有显著的调度抖动。我一般建议用 min_ms 作为主要指标avg_ms 只做参考这样可以剔除其他进程抢占 GPU 造成的偶发放大。5. cublas_benchmarks 避坑手册结果对不上、测不准、不复现的常见原因5.1 坑 1cuBLAS 默认列主序M 和 N 对调后性能完全变样现象同样的矩阵尺寸只是把 M 和 N 传反了结果 TFLOPS 掉了三分之一还不止。原因cuBLAS 的 API 遵循 Fortran 列主序约定。A 是 M×K 时存储上它是 K 个列向量拼在一起。传 M 和 K 的顺序反了相当于告诉库去读一个转置后的矩阵内存访问模式从连续变成了跨步。GPU 的 global memory 带宽对连续访问友好得多跨步访问会浪费 cache line。解决benchmark 工具里不要自作聪明去手动转置就按 API 语义传入行数和列数。如果你在循环里同时测了 M×K×N 和 K×M×N 两组它们的结果不应该视为同类数据要么分开报告要么明确标注转置状态。如果非要用行主序数据需要调用 cublasLt 的矩阵布局接口把 stride 信息显式传进去。5.2 坑 2没有 warmup第一次调用慢了十倍现象iter 设置为 100但前 3 次的耗时明显大于后面的 90 多次平均值被拉高看起来好像 cuBLAS 不稳定。原因第一次调用要触发若干初始化路径包括 cudaMalloc 的某种惰性分配、cuBLAS 内部 kernel 选择的 heuristics 加载。一次成熟的 benchmark 应该把这些一次性开销排除在统计之外。解决强制要求 warmup 参数至少大于等于 1且 warmup 跑的次数不能计入统计。如果你用的是循环注意区分预热循环和正式循环不要把两个循环的数据混在一起算。// 预热阶段 for (int i 0; i warmup; i) { cublasSgemm(handle, opA, opB, ...); } cudaDeviceSynchronize(); // 正式计时 cudaEventRecord(start); for (int i 0; i iterations; i) { cublasSgemm(handle, opA, opB, ...); } cudaEventRecord(stop); cudaDeviceSynchronize();5.3 坑 3用 clock() 计时把 CPU 等待算进去了现象测出的耗时是实际 kernel 耗时的 3 到 5 倍且每次结果波动巨大。原因cuBLAS 是异步调用cublasSgemm调用后代码立即返回kernel 实际在 GPU 上排队执行。如果用 CPU 侧的clock()或者chrono来测量 start 和 end包进来的是入队等待时间、GPU 工作队列中排队的其他 kernel 时间以及可能的 CPU-GPU 同步开销。另一种反过来的错误用cudaEventRecord但不在循环外加一次cudaDeviceSynchronize导致最后一个 kernel 还没跑完 event 的 elapsed time 是未定义的。解决只允许使用 CUDA events 做计时并且在 event 记录后调用cudaDeviceSynchronize()确保所有 kernel 完成再用cudaEventElapsedTime取时间差。严格来说更精细的测量是在循环内部每个 kernel 前后都记录一对 event取平均值但多对 event 会引入 event 本身的记录开销对只有几微秒的小 kernel 来说干扰不小。benchmark 工具里采用整体计时再除以循环次数是最稳妥的方案。5.4 坑 4TF32 默认开启精度测试被污染现象float 精度模式下新的 GPU 测出的 TFLOPS 比预期高很多但跑同一组参数算出来的结果矩阵和 CPU 参考值对不上。原因Ampere 架构以来的 cuBLAS 默认允许 Tensor Core 使用 TF32 格式处理 float 输入。TF32 的指数位和 float 一样尾数位只有 10 位精度远低于 float。它在大多数深度学习场景中表现足够好所以被设为默认但你的 benchmark 如果是纯数值精度验证这个默认行为就是个陷阱。解决在调用 benchmark 的 GEMM 之前显式设置 math modecublasSetMathMode(handle, CUBLAS_TF32_TENSOR_OP_MATH);跑基准看性能时用 TF32 没问题但跑数值对比时切换成CUBLAS_FP32_DEFAULT_MATH并且把 math mode 在输出 CSV 里记录下来。这个字段看起来琐碎实际排查精度问题时能帮你快速定位是不是被 Tensor Core 尾数截断坑了。5.5 坑 5只测一种 shape把某个结论当成普遍规律现象用 MNK4096 测试后得出「这块卡很强」的结论换到实际模型里常见的小 batch 推理场景速度掉到不可接受。原因cuBLAS 对不同维度的自适应能力很强但它不是在所有维度组合上都均匀优秀。某些尺寸组合刚好命中 kernel 的最优 tile 切分某些组合触发回退路径。只看一个 shape 的 benchmark 无法反映真实负载分布。解决benchmark 工具应该支持扫描一组维度组合而不是只跑单点。常见的做法是生成一个维度列表比如 128、256、512、1024、2048、4096 排列组合再筛掉明显无意义的组合。扫描模式下把耗时、TFLOPS 和 shape 对应关系全部输出之后再决定是看平均、看中位数还是看最差情形。真实业务中应以最接近生产 shape 的排行为准而不是拿最大的形状代表一切。6. 把基准玩出进阶价值引入回归测试与布局扫描跑通基础 benchmark 后我一般会再做两件事让工具真正变成长期资产。第一件是接入回归脚本把每次跑出的 CSV 结果存成一个基线文件之后每次升级驱动、换 CUDA 版本或改编译标志再跑一遍并做 diff。这里的关键是允许微小波动但拒绝明显下行比如阈值设为 5%。如果同一组 shape 的 TFLOPS 掉超过 5%基本可以确定是哪次改动引入了性能回退。第二件事是矩阵布局扫描。除了标准连续的列主序布局我在实际项目里遇到过多次需要非连续布局的场景比如 A 只有 m 行中的一部分参与计算或者 B 是按子块临时拷贝出来的此时 stride 不等于行数和列数。cublasLt 允许显式定义 lda、ldb、ldc 的步长benchmark 工具如果只测连续布局就会在这些真实场景里失去预测力。cublasLtMatrixLayout_t layoutA, layoutB, layoutC; cublasLtMatrixLayoutCreate(layoutA, CUDA_R_32F, m, k, lda); cublasLtMatrixLayoutCreate(layoutB, CUDA_R_32F, k, n, ldb); cublasLtMatrixLayoutCreate(layoutC, CUDA_R_32F, m, n, ldc); cublasLtMatmul(ltHandle, alpha, ...);这段代码的价值在于你可以把 lda 设置在 m 之上比如 m16模拟出真实对齐场景的 padding。有次某个部署环境的输入张量为了对齐申请了 528 字节每行的缓冲比连续布局多跑出 8% 的带宽损失用布局扫描后才发现原因。我自己的教训是基准工具最容易出错的不是实现而是参数和模式没被记录就发布了结果。但凡输出里缺一个 dtype 或 math mode 字段后续对比就是浪费时间。现在我的习惯是每次跑完都让工具自动把环境信息、版本号、mode 和维度全部打进 CSV再谈分析。单元测试和基线都是虚的可复现的记录才是真本钱。希望这个 benchmark 工具能帮你少走一段计时器的弯路早日把矩阵乘法的性能边界摸透。本文还有配套的精品资源点击获取