简介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×klda 最小取 m如果 transa 是CUBLAS_OP_TA 的逻辑形状变成 k×m但内存里还是原来的排列lda 的取值会变。基准测试里最常用的组合是 N/T也就是 A 不转置、B 转置这个组合在神经网络推理场景最常见测试时最好把四种转置组合全覆盖。GEMV 的参数就少很多核心是cublasSgemv(handle, trans, m, n, alpha, A, lda, x, incx, beta, y, incy)。GEMV 是矩阵乘向量A 的形状是 m×nx 是 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 读入 Am*n 元素、xn 元素、ym 元素写出 ym 元素按 FP32 算总字节数约等于4 * (m*n n 2*m)用这个去除耗时得到 GB/s。shape 矩阵怎么设计决定了基准有没有参考价值。我常用的方案是三类 shapeShape 类型典型范围主要目的算力上限型m/n/k 均为 4096~8192测峰值 TFLOPS中等维度型m/n/k 在 512~2048贴近视觉模型卷积降维后的矩阵乘小 shape 型m/n/k 在 16~256测 kernel 启动开销和切换点这三类要分开跑不能混在一起取平均值。混在一起的问题很明显大 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::vectorfloat hA(m * k, 1.0f); std::vectorfloat hB(k * n, 1.0f); std::vectorfloat 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_TB 的逻辑列数变成 n但 ldb 依然取 k这个细节很多人会在改转置时搞错。参数说明repeat是重复次数后面计时部分会用到建议至少 100 次起步alpha1.0, beta0.0是最常见配置但如果你要测类似残差连接那样的beta1.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 正确性校验没有校验的基准测出来的是噪声基准测试的另一个经典翻车点是只看性能不看正确性。某次我调一个 m2048 的 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::vectorfloat 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、带宽和校验状态。不要输出中间过程的调试信息否则几千行日志里找一条记录会让人崩溃。输出列含义读法shape_id输入文件行号定位 shapem/n/k输入维度必须回显防止文件顺序改错time_ms单次平均耗时看绝对耗时是否合理tflops计算型指标对比峰值占比gbps访存型指标GEMV 主要看这个statusok / err校验失败直接标 err读结果时有个经验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 变化的曲线而不是只看表格。我一般固定 mnk从 64 扫到 8192 做对数扫描观察耗时曲线里的拐点。拐点位置往往对应 cuBLAS 从朴素 kernel 切到分块 kernel 或者切到 Tensor Core kernel 的位置。找到这些切换点后你在做算子融合或者选择手动 kernel 时就有了依据哪个 shape 区间值得投入手动优化哪个区间直接用 cuBLAS 就够了。这个做法在业务里很实用。某次我在做某图像处理 Demo 的矩阵乘优化发现 m512、n2048、k4096 的 shape 落在 cuBLAS 的切换边界上耗时是相邻 shape 的近两倍。后来就是在基准里把 k 维度微调成 4352 并配合 padding避开了切换点整体推理延迟降了 30%。没有基准曲线这种收益根本无从下手。6.2 基准做回归把关键 shape 固化进日常检查基准不只可以用在人肉决策也可以固化成一个回归工具。我会把业务中最常用的 10 个 shape 单独拉出来做成一个快速测试集每次切换 cuBLAS 版本、驱动版本或者 GPU 环境后跑一遍把 TFLOPS 变化量打印出来。如果某个 shape 的 TFLOPS 掉超过 5%就说明环境变化影响了这条路径需要定位。这个做法比全量跑几百个 shape 高效得多适合放在日常开发流程里持续执行。6.3 最后补一句我的习惯我每次跑完基准不会立刻删掉跑批脚本而是把 shapes.txt 和一个简单的解析脚本留在项目里下次有新卡或者新库版本直接复用。另一个习惯是永远把校验和计时写在同一份代码里哪怕只是临时跑一个 shape。为省这几行代码吞掉一个假数据后续排查的代价会更大。这套框架不算复杂但足够撑起从选型评估到回归检查的大部分需求。希望帮到你。本文还有配套的精品资源点击获取