企业官网建设流程全解析

简介：cublas_benchmarks是一套面向CUDA开发者与性能优化场景的基准测试工程，专注于测量CuBLAS库中GEMM（通用矩阵乘法）与GEMV（通用矩阵向量乘法）在不同矩阵规模下的实际执行时间，为科学计算、机器学习和深度学习中的密集线性代数运算提供调优参考。压缩包共10个文件，包括4个头文件、4个CUDA源码、1个makefile和1个readme，整体仅13KB，结构紧凑，便于快速编译与二次修改。项目通过初始化数据、调用CuBLAS函数、同步GPU并记录耗时的方式输出性能结果，使用者可借此分析特定硬件配置下的计算瓶颈，并对比不同CUDA版本或优化策略的差异。该资源已有1095人学习下载，对于需要自行搭建矩阵运算性能测试环境、评估GPU加速效果的开发者来说，是一份轻量且可直接运行的实践参考。

1. 为什么我会专门写一份测 GEMM/GEMV 的基准：先想清楚要测什么

很多人以为矩阵乘性能测试就是把 cuBLAS 的 sgemm 丢进循环里跑一遍，然后除以墙钟时间。真正上手做算子选型或者评估某块 GPU 的算力天花板时会发现，同一组 M/N/K，换一个 lda、换一种转置组合、甚至只是把重复次数从 10 改成 100，得出的 TFLOPS 能差出一倍。这不是玄学，是 cuBLAS 内部 kernel 选择逻辑、内存布局与计时方式叠加出来的结果。

cublas_benchmarks 这类基准要解决的就是这件事：把 GEMM 和 GEMV 的调用用统一框架包一层，按固定 shape 列表跑批，给出稳定的耗时和带宽读数。它适合三类人：做算子选型的人、评估新环境算力的人、以及想摸清某张卡在不同 shape 下性能拐点的人。我最初自己手写这个框架，是因为某公司内部做某跨平台系统的算子优化，需要给几十组 shape 做前后对比，而现成工具要么只测固定矩阵、要么没法控制精度和转置组合。这篇文章就把整套做法讲清楚，从接口参数到计时框架，再到最容易翻车的地方。

2. 基准在测什么：GEMM、GEMV 的变体、转置与五个必填参数

2.1 从矩阵乘公式到 cuBLAS 接口：参数和内存布局是核心

要写基准，第一步是把 cuBLAS 的函数签名看透。GEMM 的数学定义是：

C = alpha * op(A) * op(B) + beta * C

其中 op 可以是转置或非转置。cuBLAS 的 sgemm 接口长这样：

cublasStatus_t cublasSgemm( cublasHandle_t handle, cublasOperation_t transa, // 对 A 是否转置 cublasOperation_t transb, // 对 B 是否转置 int m, int n, int k, // 三个维度 const float* alpha, const float* A, int lda, // A 的 leading dimension const float* B, int ldb, const float* beta, float* C, int ldc );

这里最容易让新手困惑的其实是两个点。第一，cuBLAS 是列主序存储，也就是说矩阵是按列优先连续存放在内存里的；如果你平时在 C++ 里用A[i][j]写代码，那是在按行主序处理，直接把指针传给 cuBLAS 时 lda 要特别小心，否则生产出来的精度数值没问题，但实际计算的矩阵和你想的完全不是同一个。第二，m/n/k 与 op 的操作是有对应关系的：如果 transa 是CUBLAS_OP_N，那么 A 的形状是 m×k，lda 最小取 m；如果 transa 是CUBLAS_OP_T，A 的逻辑形状变成 k×m，但内存里还是原来的排列，lda 的取值会变。基准测试里最常用的组合是 N/T，也就是 A 不转置、B 转置，这个组合在神经网络推理场景最常见，测试时最好把四种转置组合全覆盖。

GEMV 的参数就少很多，核心是cublasSgemv(handle, trans, m, n, alpha, A, lda, x, incx, beta, y, incy)。GEMV 是矩阵乘向量，A 的形状是 m×n，x 是 n 维，y 是 m 维。别看它简单，性能模型完全不同：GEMM 是计算密集，GEMV 是访存密集，后者的耗时基本由向量和矩阵的带宽决定，所以测试指标要用有效带宽而不是 TFLOPS。

2.2 性能指标怎么算：TFLOPS、有效带宽与 shape 矩阵的设计

基准测试的产出必须是可以横向比较的指标，而不是裸的毫秒数。GEMM 的计算量是 2 * m * n * k 次浮点运算，乘 2 是因为乘法和加法各算一次，所以 FLOPS 就是2.0 * m * n * k / time，再除以 1e12 就是 TFLOPS。GEMV 我一般不看 FLOPS，因为矩阵向量乘在绝大多数 GPU 上都跑不满计算单元，直接用有效带宽：一次 GEMV 读入 A（m*n 元素）、x（n 元素）、y（m 元素），写出 y（m 元素），按 FP32 算总字节数约等于4 * (m*n + n + 2*m)，用这个去除耗时得到 GB/s。

shape 矩阵怎么设计，决定了基准有没有参考价值。我常用的方案是三类 shape：

这三类要分开跑，不能混在一起取平均值。混在一起的问题很明显：大 shape 的耗时可能是小 shape 的几百倍，平均下来小 shape 的性能信息完全被淹没，什么结论都读不出来。真正的做法是每一组 shape 单独计时、单独记录，最后以 shape 为行输出一张表。

另外要注意数据类型。FP32 的 sgemm 和 FP16 的 hgemm 走的是完全不同的 kernel 路径。如果标题只写 GEMM/GEMV 性能，起步阶段把 FP32 跑扎实就够；但你的基准如果面向深度学习场景，最好从一开始就把精度列档案设计进去，因为后续几乎一定会补 FP16。

3. 搭一个能跑出数字的最小基准：计时、warmup 与正确性校验

3.1 打开 cuBLAS 句柄与参数解析：一次初始化全局复用

cuBLAS 的调用必须先创建一个句柄，这个句柄负责维护设备端的工作空间和调度的上下文。句柄创建后不要反复创建销毁，基准测试里这是常见性能噪音来源。以下你用 C++ 搭一个最小骨架：

#include <cuda_runtime.h> #include <cublas_v2.h> #include <vector> #include <string> #include <chrono> int main(int argc, char* argv[]) { int m = std::stoi(argv[1]); int n = std::stoi(argv[2]); int k = std::stoi(argv[3]); int repeat = std::stoi(argv[4]); cublasHandle_t handle; cublasCreate(&handle); // 全局只创建一次 // 设备端数据准备 std::vector<float> hA(m * k, 1.0f); std::vector<float> hB(k * n, 1.0f); std::vector<float> hC(m * n, 0.0f); float *dA, *dB, *dC; cudaMalloc(&dA, m * k * sizeof(float)); cudaMalloc(&dB, k * n * sizeof(float)); cudaMalloc(&dC, m * n * sizeof(float)); cudaMemcpy(dA, hA.data(), m * k * sizeof(float), cudaMemcpyHostToDevice); cudaMemcpy(dB, hB.data(), k * n * sizeof(float), cudaMemcpyHostToDevice); cudaMemcpy(dC, hC.data(), m * n * sizeof(float), cudaMemcpyHostToDevice); // 测试参数 float alpha = 1.0f, beta = 0.0f; cublasOperation_t transa = CUBLAS_OP_N; cublasOperation_t transb = CUBLAS_OP_N; int lda = m, ldb = k, ldc = m; // 调用入口 for (int i = 0; i < repeat; i++) { cublasSgemm(handle, transa, transb, m, n, k, &alpha, dA, lda, dB, ldb, &beta, dC, ldc); } // ... 后续计时与打印 cublasDestroy(handle); }

逻辑说明：这个骨架把 host 端数据和 device 端数据分离，host 数据只用于初始化，真正的计算全部在显存里。lda = m对应的是列主序下非转置 A 的最小合法值；如果你后续把 transb 改成CUBLAS_OP_T，B 的逻辑列数变成 n，但 ldb 依然取 k，这个细节很多人会在改转置时搞错。

参数说明：repeat是重复次数，后面计时部分会用到，建议至少 100 次起步；alpha=1.0, beta=0.0是最常见配置，但如果你要测类似残差连接那样的beta=1.0场景，需要单独跑一组，因为 beta 的不同取值会影响 cuBLAS 是否走融合 kernel 分支。

3.2 计时封装：用 cudaEvent 而不是墙钟，加 warmup 再加中位数

计时是整个基准里最影响可信度的环节。std::chrono在 kernel 异步执行下完全不可用——CPU 计时器读到的只是提交时间，不是 GPU 完成时间。正确做法是 cudaEvent，它能记录 GPU 时间轴上的事件。

void timed_gemm(cublasHandle_t handle, cublasOperation_t transa, cublasOperation_t transb, int m, int n, int k, float alpha, const float* dA, int lda, const float* dB, int ldb, float beta, float* dC, int ldc, int repeat, float& ms) { // warmup：让 cuBLAS 完成 kernel 选择、cuBLASLt 启发式搜索、各种初始化 cublasSgemm(handle, transa, transb, m, n, k, &alpha, dA, lda, dB, ldb, &beta, dC, ldc); cudaDeviceSynchronize(); cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop); // 单独记录一组时间：事件尽量贴近 kernel，避免 memcpy 混入 cudaEventRecord(start); for (int i = 0; i < repeat; i++) { cublasSgemm(handle, transa, transb, m, n, k, &alpha, dA, lda, dB, ldb, &beta, dC, ldc); } cudaEventRecord(stop); cudaEventSynchronize(stop); cudaEventElapsedTime(&ms, start, stop); ms /= repeat; // 多次计时取中位数，降低系统降频影响 cudaEventDestroy(start); cudaEventDestroy(stop); }

逻辑说明：warmup 那一次调用非常重要。cuBLAS 在第一次拿到陌生 shape 时要做内部 kernel 筛选，某些算法甚至要跑一遍 profile 才能选最优，这段时间可能比正式计算还长，不排除掉的话你的结果会有一个固定的大幅偏移。cudaEvent 记录的时间范围只包住 GPU 上的 kernel 执行，cudaEventSynchronize(stop)会阻塞 CPU 直到 GPU 执行到 stop 事件，保证读到的 ms 是完整的。

注意 repeat 的取值与耗时成反比。对 4096 级别的大 shape，一次 kernel 可能就是几毫秒，repeat 取 50 就够；对 64 级别的小 shape，一次可能只有几十微秒，repeat 至少给 500，否则事件分辨率不够。我会在参数里做一个自适应：预计单次耗时小于 0.1ms 时自动把 repeat 翻倍，这是用时间换稳定性。

3.3 正确性校验：没有校验的基准测出来的是噪声

基准测试的另一个经典翻车点是只看性能不看正确性。某次我调一个 m=2048 的 sgemm，发现测出的 TFLOPS 比理论峰值还高，细查才发现是 lda 传错了，cuBLAS 实际读取的矩阵只有真实数据的一小部分，计算量天然就小了一半，性能数字自然虚高。所以基准框架里必须有校验环节：

void verify_gemm(cublasHandle_t handle, int m, int n, int k, const float* dA, int lda, const float* dB, int ldb, float* dC, int ldc) { // 用 CPU 做一个朴素三层循环参考实现 std::vector<float> hC(m * n, 0.0f); // 注意这里 hA hB 从设备拷回 host 的代码省略，读者可自行补全 for (int i = 0; i < m; i++) { for (int j = 0; j < n; j++) { float sum = 0.0f; for (int p = 0; p < k; p++) { sum += hA[i * k + p] * hB[p * n + j]; } hC[i * n + j] = sum; } } float max_err = 0.0f; float* dC_ref; cudaMalloc(&dC_ref, m * n * sizeof(float)); cudaMemcpy(dC_ref, hC.data(), m * n * sizeof(float), cudaMemcpyHostToDevice); // 计算最大绝对误差 // 阈值一般放 1e-3 到 1e-2 之间，FP32 累加误差对 2048 维已经可能到 1e-2 float tolerance = 1e-2f; // 超过 tolerance 则打印告警并标记该 shape 结果不可信 }

逻辑说明：校验的阈值要按矩阵规模放宽，不能用单精度 epsilon。K 维度越大，浮点累加误差越大，我一般按tolerance = 1e-2f起设，如果校验失败会直接把时间和 TFLOPS 标成无效数据，而不是让一条假数据混进结果表。这个环节看起来费事，但它救回过我好几次因为转置组合调整引入的错误配置。

4. 如何组织跑批：shape 列表、精度组合与结果输出格式

4.1 用一个 shape 列表驱动整晚测试

单个 shape 的手动测试只适合 debug，真正做基准要批量跑。我的做法是把 shape 列表写在一个独立文件里，跑批程序读取它、循环执行。这样后续加 shape 不需要改代码，而且跑批时每个 shape 是独立进程还是共享进程可控。

# shapes.txt 示例：每行是 m n k transa transb 2048 2048 2048 N N 4096 4096 4096 N T 8192 1024 4096 N N 128 4096 4096 N N 64 64 64 N N 256 256 512 T N 512 1024 2048 T T

主程序只需要读入这几列，然后循环调用timed_gemm。跑批顺序上有个实用技巧：把所有 shape 按计算量从小到大排序，这样启动阶段不会让几个超大 shape 霸占 GPU 太久，遇到超时还能及时跳过。另一个技巧是每个 shape 之间加一次cudaDeviceSynchronize()，防止后续 shape 的数据拷贝与前面的 kernel 排队，污染时间统计。

为什么用文件而不是命令行参数？命令行参数适合跑单次实验，但你做 shape 扫描时经常要跑几十上百组，且需要回到同一组 shape 做库版本对比。文件驱动意味着你换一个 cuBLAS 版本后，只要改脚本里的库路径，shape 集合完全复用，对比结果才有意义。

4.2 输出结果怎么读：区分 compute-bound 与 memory-bound

跑批的输出格式我见过各种设计，但最核心的原则是：一行一个 shape，列包含耗时、TFLOPS、带宽和校验状态。不要输出中间过程的调试信息，否则几千行日志里找一条记录会让人崩溃。

读结果时有个经验：GEMM 看 TFLOPS 对硬件峰值的占比，如果某组 shape 下达到 60% 以上说明 kernel 选得不错；如果只有 10%，先别怀疑硬件，优先排查是不是 shape 太小、机器没进最大频率或者 lda 设置导致访存模式碎片化。GEMV 看带宽，带宽曲线通常是先随 n 增大快速上升，到某个点后趋平，趋平前的点往往是计算启动开销主导。

4.3 精度与转置组合怎么纳入测试矩阵

只测 FP32 N/N 组合只能回答"sgemm 性能如何"，不能回答业务问题。我更推荐在跑批时把精度和转置作为组合维度：FP32 N/N、FP32 N/T、FP32 T/N、FP16 N/N。这样一个 shape 最多产生四行结果，但这才贴近真实场景。比如卷积网络的 im2col 之后往往是 N/T 组合，Transformer 的 attention 打分是 T/N 组合，只测单一组合会导致你基于基准做的选型在真实网络里偏差很大。

精度组合的注意点是 API 要切换：

// FP16 GEMM 不能再用 cublasSgemm，要用支持混合精度的入口 cublasGemmEx(handle, transa, transb, m, n, k, &alpha_one, dA_half, CUDA_R_16F, lda, dB_half, CUDA_R_16F, ldb, &beta_zero, dC_half, CUDA_R_16F, ldc, CUDA_R_32F, // 计算精度，FP16 乘积累加到 FP32 再写回 CUBLAS_GEMM_DEFAULT_TENSOR_OP_ALGO);

逻辑说明：cublasGemmEx比cublasSgemm多出数据类型和算法选择两个维度。计算精度选CUDA_R_32F是精度优先的配置，选CUDA_R_16F是纯粹走 Tensor Core 的配置，性能差不少，但误差也会变大。基准里两种都要测，因为实际工程里需要自行权衡。另外，转置组合会影响能否用 Tensor Core 的特定指令布局，这也是为什么建议把所有转置组合都跑进结果，而不是默认只测 N/N。

5. 避坑：写 GEMM/GEMV 基准最容易翻车的五个点

5.1 性能数字正常，结果却完全算错：lda 与列主序的坑

现象：测出的 TFLOPS 稳定在 70% 峰值以上，看起来一切正常；某次顺手加了校验，发现最大误差超过 10。

原因：宿主代码按行主序初始化矩阵，lda 直接填了 m，但 cuBLAS 按列主序解释数据，相当于实际矩阵被扭转了 90 度，内存读取区间不完整。性能看起来不差是因为 kernel 计算量本身没变，只是算的内容是错的。

解决：初始化数据时就用列主序索引公式A[i + j * lda]，或者干脆把 hA/hB 的生成函数改成列主序标准。校验环节必须保留，这是我上面花一整节强调校验的原因。

5.2 第一次调用奇慢，后续变快：cuBLAS 的 kernel autotuning

现象：repeat 100 次，前 2 次每次 8ms，后面稳定在 3ms，平均值被抬高 10% 以上。

原因：cuBLAS 针对陌生 shape 会运行内部启发式搜索，某些算法选项会做实际执行来评估代价。

解决：正式计时前先跑一次同 shape 的调用做 warmup，然后单独计时。注意 warmup 的调用必须和正式调用用同一个 handle，因为 autotuning 的缓存绑定在 handle 上，换个 handle 等于重新来一遍。

5.3 GPU 没跑到满频：功耗与频率爬升被算进时间里

现象：连续跑大 shape 时，前几轮 TFLOPS 偏低，后面逐渐升高到稳定值；如果休息几秒再跑，又会出现短暂回落。

原因：GPU 频率从 idle 状态爬升需要时间，基准跑批如果每个 shape 之间间隔太久，频率掉下去，导致数字偏低。

解决：在正式测试前加一段大 shape 预热循环让 GPU 进入稳定温度与功耗状态。跑批过程中不要用 sleep 隔开 shape，连续执行保持热态；如果一定要间隔，至少先把 GPU 跑热再测。

5.4 CPU 端 memcpy 混入计时：带宽数字假高

现象：GEMV 测出的有效带宽超过设备理论显存带宽上限，明显不合理。

原因：有些封装会把cudaMemcpyAsync放进计时区间，或者 host 端传参没做异步处理，CPU 在 GPU kernel 执行期间偷跑，事件记录的时间比实际 kernel 短。

解决：所有数据拷贝和 shape 解析都放在事件记录之前，计时区间内只放 kernel 调用。把cudaDeviceSynchronize放在事件之外，避免 CPU 等待开销污染 GPU 时间。

5.5 不同转置组合测出完全不同的性能，误判为 bug

现象：N/N 和 N/T 的 TFLOPS 相差 30% 以上，怀疑 cuBLAS 对转置有性能缺陷。

原因：cuBLAS 对不同内存布局会选不同 kernel，部分 kernel 对不连续的 B 访问效率低，但这种差异是真实存在的，不是 bug。N/T 在某些 shape 下慢，可能正对应业务代码里该转置的组合，需要另寻对策。

解决：不要试图让所有转置组合都一样快，而是把基准结果当作真实参考。后续优化如果必须换布局，直接同时改数据排布和 lda，再跑基准看收益，而不是相信直觉。

6. 进阶：把基准结果做成选型决策而不是一张性能表

6.1 用 shape 扫描找 kernel 切换点

当你的基准已经把几十个 shape 跑完，下一步要做的是画出耗时随 shape 变化的曲线，而不是只看表格。我一般固定 m=n=k，从 64 扫到 8192 做对数扫描，观察耗时曲线里的拐点。拐点位置往往对应 cuBLAS 从朴素 kernel 切到分块 kernel 或者切到 Tensor Core kernel 的位置。找到这些切换点后，你在做算子融合或者选择手动 kernel 时就有了依据：哪个 shape 区间值得投入手动优化，哪个区间直接用 cuBLAS 就够了。

这个做法在业务里很实用。某次我在做某图像处理 Demo 的矩阵乘优化，发现 m=512、n=2048、k=4096 的 shape 落在 cuBLAS 的切换边界上，耗时是相邻 shape 的近两倍。后来就是在基准里把 k 维度微调成 4352 并配合 padding，避开了切换点，整体推理延迟降了 30%。没有基准曲线，这种收益根本无从下手。

6.2 基准做回归：把关键 shape 固化进日常检查

基准不只可以用在人肉决策，也可以固化成一个回归工具。我会把业务中最常用的 10 个 shape 单独拉出来，做成一个快速测试集，每次切换 cuBLAS 版本、驱动版本或者 GPU 环境后跑一遍，把 TFLOPS 变化量打印出来。如果某个 shape 的 TFLOPS 掉超过 5%，就说明环境变化影响了这条路径，需要定位。这个做法比全量跑几百个 shape 高效得多，适合放在日常开发流程里持续执行。

6.3 最后补一句我的习惯

我每次跑完基准不会立刻删掉跑批脚本，而是把 shapes.txt 和一个简单的解析脚本留在项目里，下次有新卡或者新库版本直接复用。另一个习惯是永远把校验和计时写在同一份代码里，哪怕只是临时跑一个 shape。为省这几行代码吞掉一个假数据，后续排查的代价会更大。这套框架不算复杂，但足够撑起从选型评估到回归检查的大部分需求。希望帮到你。

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