企业官网建设流程全解析

简介：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×K，B 是 K×N，C 是 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 * y，A 是 M×N，x 是 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 * 4（float 类型），再算上写回 y 向量的 M * 4 字节，总共大约是（M * N + M）* 4 字节，用 字节数 / 时间 / 1e9 得到 GB/s。

另外有一个容易忽略的指标：与理论峰值的百分比。拿 GEMM 的实测 TFLOPS 除以显卡的理论 TFLOPS（float 精度下），得到一个效率百分比。业界经验是方形大矩阵能跑到 85% 以上就算正常，明显低于这个值就要检查是不是没关 ECC、驱动降频或散热问题。理论峰值可以从设备属性里拿到 SM 数量和时钟频率大致估算，不必精读汇编指令。

3. 把 cublas_benchmarks 跑起来：依赖、编译与最小命令行

3.1 环境检查与编译流程

常见的做法是先确认两件事：CUDA 工具链版本和显卡算力。cuBLAS 是随 CUDA toolkit 一起发布的动态库，如果你的程序用了 cuBLAS，运行时还需要确保动态库路径正确。执行下面的命令做前置检查：

nvcc --version nvidia-smi --query-gpu=name,compute_cap,memory.total --format=csv

第一行看编译器版本，第二行看驱动识别的显卡名、算力代数和显存总量。算力代数是 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_TYPE=Release 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: M=4096 N=4096 K=4096 alpha=1.00 beta=0.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 与数据类型

下表整理了关键参数的含义和建议初始值，这个表已经够覆盖绝大多数使用场景。

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 倍以上。拿模拟数据举例：在一张中等定位的显卡上，M=N=K=4096 可能跑到 20+ TFLOPS，而 M=1024，N=64，K=4096 可能只有 4 TFLOPS 左右。原因有两个：一是瘦长矩阵的切分被迫使用更小的 tile，无法把 Tensor Core 的 16×16×16 计算块打满；二是很小的 N 会让 C 矩阵的写回访存次数按 tile 放大，访存占比上升。

GEMV 的差异更多来自矩阵的行数和列数比例。A 矩阵按行主序存储时，行数大意味着 y 向量的写回次数多，写回可能遇到非连续地址。某个测试里 M=1,000,000，N=16 的 GEMV 跑到的带宽不到 M=100,000，N=256 的一半，就是因为 y 向量的写回规模和缓存命中率都变了。

4.3 一个可信的基准输出应该长什么样

为了让结果能进报告或做回归对比，工具的输出别只打印正文里的那种格式化文本，还应该输出 CSV 行。这样同类数据可以直接叠加到同一个表格里。推荐的字段有：时间戳、GPU 名称、CUDA 版本、op_type（GEMM 或 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 坑 1：cuBLAS 默认列主序，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 坑 4：TF32 默认开启，精度测试被污染

现象：float 精度模式下，新的 GPU 测出的 TFLOPS 比预期高很多，但跑同一组参数算出来的结果矩阵和 CPU 参考值对不上。

原因：Ampere 架构以来的 cuBLAS 默认允许 Tensor Core 使用 TF32 格式处理 float 输入。TF32 的指数位和 float 一样，尾数位只有 10 位，精度远低于 float。它在大多数深度学习场景中表现足够好，所以被设为默认，但你的 benchmark 如果是纯数值精度验证，这个默认行为就是个陷阱。

解决：在调用 benchmark 的 GEMM 之前，显式设置 math mode：

cublasSetMathMode(handle, CUBLAS_TF32_TENSOR_OP_MATH);

跑基准看性能时用 TF32 没问题，但跑数值对比时切换成CUBLAS_FP32_DEFAULT_MATH，并且把 math mode 在输出 CSV 里记录下来。这个字段看起来琐碎，实际排查精度问题时能帮你快速定位是不是被 Tensor Core 尾数截断坑了。

5.5 坑 5：只测一种 shape，把某个结论当成普遍规律

现象：用 M=N=K=4096 测试后，得出「这块卡很强」的结论，换到实际模型里常见的小 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 之上比如 m+16，模拟出真实对齐场景的 padding。有次某个部署环境的输入张量为了对齐申请了 528 字节每行的缓冲，比连续布局多跑出 8% 的带宽损失，用布局扫描后才发现原因。

我自己的教训是：基准工具最容易出错的不是实现，而是参数和模式没被记录就发布了结果。但凡输出里缺一个 dtype 或 math mode 字段，后续对比就是浪费时间。现在我的习惯是每次跑完都让工具自动把环境信息、版本号、mode 和维度全部打进 CSV，再谈分析。单元测试和基线都是虚的，可复现的记录才是真本钱。希望这个 benchmark 工具能帮你少走一段计时器的弯路，早日把矩阵乘法的性能边界摸透。

本文还有配套的精品资源，点击获取