简介cublas_benchmarks 是一份面向 CUDA/C 开发者与高性能计算学习者的 CuBLAS 性能测试工程围绕 GEMM、GEMV 两类基础线性代数算子在 GPU 上的执行表现展开可用于科学计算、机器学习与深度学习场景下的性能评估和调优参考。压缩包共 10 个文件约 13KB以 4 个 .cu CUDA 源文件、4 个 .h 头文件为主并配有 Makefile 与 README分别承载测试逻辑实现、数据结构与函数声明、编译构建和使用说明目录组织简洁便于按功能模块阅读与二次扩展。已有 1095 人学习下载。代码通过设置不同矩阵与向量规模调用 CuBLAS 函数配合 GPU 同步与耗时记录帮助读者观察不同参数、CUDA 版本和硬件配置下的性能差异理解基准测试的完整流程并为进一步优化计算策略、选择合适算子实现提供可复用参考对需要复现实验、对比实现或排查性能瓶颈的读者尤具价值。1. 同一张卡跑出三倍差距cuBLAS 的 GEMM 与 GEMV 基准到底在测什么同一张卡、同一组方阵两份 cublas_benchmarks 跑出来的 GEMM 有效算力能差出两三倍差别往往不在硬件而在计时方式、L2 缓存命中情况和参数传法。这类基准做的事很具体把 cuBLAS 里的 GEMM矩阵乘矩阵和 GEMV矩阵乘向量按不同维度、不同精度、不同转置组合各跑一遍记录耗时并换算成算力TFLOPS和带宽GB/s最终得到一条“这张卡在当前形状下能跑多快”的曲线。它对应两类高频需求自研 kernel 到底比 cuBLAS 快还是慢以及硬件选型时瓶颈在算力侧还是显存带宽侧。写自定义算子要基线、调推理吞吐要找瓶颈、做验收要一个可复现数字的人都适合按这套流程跑一遍但必须先把测量规则定对。2. 搭起基准骨架依赖、目录与一个不误导人的计时函数基准代码本身不难难的是让它报出来的数字可被信任。我见过太多版本是用std::chrono包住cublasSgemm调用然后再cudaDeviceSynchronize()看起来合理实际把主机端提交耗时也算进去了矩阵小的时候误差能到百分之几十。所以整个骨架要从计时函数开始搭而不是从 kernel 开始搭。2.1 依赖清单与目录结构需要的东西不多但每一样都要能自证存在避免写到一半才发现链接不上。组件作用自检方式CUDA Toolkitnvcc cuBLAS 运行库编译与链接基准程序nvcc --version确认$CUDA_HOME/lib64下有libcublas.so一块可用的 GPU 与匹配驱动实际执行跑 SDK 自带的 deviceQuery 样例看cudaGetDeviceProperties是否正常返回C 编译器宿主侧代码g --version与 nvcc 的宿主编译器要求一致即可Python 3 matplotlib把结果画成曲线仅在出图时用到不影响测量命令行 profiler可选定位是计算受限还是访存受限能挂上就能看 SM 利用率与 DRAM 吞吐目录按“源码 / 脚本 / 结果”三分结果目录一律带时间戳避免两次实验互相覆盖cublas_benchmarks/ ├── src/ │ ├── timer.cuh # 计时与错误检查宏 │ ├── bench_gemm.cu # GEMM 扫描 │ └── bench_gemv.cu # GEMV 扫描 ├── scripts/ │ ├── run_sweep.sh # 批量跑不同形状 │ └── plot.py # 读 csv 出曲线 └── results/ └── 20240101_1200/ ├── gemm.csv └── gemv.csv目录本身不影响性能但它决定了你能不能在一个月后复现当时那条曲线。血泪经验是形状、精度、迭代次数、flush 开关这些参数必须一起落进 csv只存耗时的话回头根本解释不了那个尖峰是怎么来的。2.2 用 cudaEvent 写一个能复用的计时函数核心思路是把“测什么”和“怎么测”分开被测体是一个可调用对象计时器只负责 warmup、迭代、事件打点以及把每次迭代前的缓存清理动作单独计价再扣掉。// src/timer.cuh #pragma once #include cuda_runtime.h #include cublas_v2.h #include cstdio #include cstdlib #include functional #define CK(x) do { cudaError_t _e (x); \ if (_e ! cudaSuccess) { std::fprintf(stderr, CUDA %s %s:%d\n, \ cudaGetErrorString(_e), __FILE__, __LINE__); std::exit(1); } } while (0) #define CB(x) do { cublasStatus_t _s (x); \ if (_s ! CUBLAS_STATUS_SUCCESS) { std::fprintf(stderr, cuBLAS %d %s:%d\n, \ (int)_s, __FILE__, __LINE__); std::exit(1); } } while (0) // body: 被测调用pre: 每次迭代前的 L2 清理不做就传空 lambda inline float timeit(const std::functionvoid() body, const std::functionvoid() pre, int warmup, int iters) { for (int i 0; i warmup; i) { pre(); body(); } CK(cudaDeviceSynchronize()); // warmup 结果全部丢弃 cudaEvent_t e0, e1; CK(cudaEventCreate(e0)); CK(cudaEventCreate(e1)); // 先单独量一遍 pre 的成本稍后从总时间里扣掉 CK(cudaEventRecord(e0)); for (int i 0; i iters; i) pre(); CK(cudaEventRecord(e1)); CK(cudaEventSynchronize(e1)); float pre_ms 0.f; CK(cudaEventElapsedTime(pre_ms, e0, e1)); CK(cudaEventRecord(e0)); for (int i 0; i iters; i) { pre(); body(); } CK(cudaEventRecord(e1)); CK(cudaEventSynchronize(e1)); float all_ms 0.f; CK(cudaEventElapsedTime(all_ms, e0, e1)); cudaEventDestroy(e0); cudaEventDestroy(e1); return (all_ms - pre_ms) / iters; // 单次迭代毫秒数 }逻辑说明cudaEvent打在流上记录的是设备侧时间点天然排除了主机提交开销这是它比std::chrono靠谱的根本原因。用std::function包被测体是为了让同一个计时器同时服务 GEMM、GEMV、batched 三种调用不必写三份。参数说明warmup建议不低于 5cuBLAS 句柄在首次调用时要做内部初始化同时 GPU 时钟也需要一段时间的负载才能稳定频率这两件事都会让第一次迭代明显偏慢iters不能照抄应该按“单次耗时 × 迭代数 ≥ 200 ms”反推否则事件分辨率和时钟抖动会主导结果pre的扣减是必须的因为清理缓存的 kernel 与 GEMM 在同一个流上串行执行会直接计入总量。清理缓存的实现不必复杂写一遍就行关键是容量要对齐设备属性// 按设备的 L2 实际容量分配不要写死一个数字 cudaDeviceProp prop; CK(cudaGetDeviceProperties(prop, 0)); size_t flush_bytes (size_t)prop.l2CacheSize * 2; float* d_flush; CK(cudaMalloc(d_flush, flush_bytes)); auto flush []{ CK(cudaMemsetAsync(d_flush, 0, flush_bytes)); };提示l2CacheSize在不同架构上差别很大写死 40 MB 或 60 MB 都会在某类卡上失效。翻车方式很隐蔽——缓存没挤干净测出来的算力反而“超过峰值”还以为是代码写得好。2.3 编译参数与第一次 baseline 跑通编译只需要-lcublas但几个开关值得固定下来nvcc -O3 -stdc14 -archnative -lineinfo \ src/bench_gemm.cu -lcublas -o build/bench_gemm ./build/bench_gemm --m 4096 --n 4096 --k 4096 --dtype fp32 --iters 50-O3影响的是宿主侧循环对 device 代码影响有限但没坏处-archnative让编译器按当前设备生成代码省得手填架构号填错导致跑在兼容模式上-lineinfo保留行号信息配合 profiler 能直接定位到具体调用点。命令行参数我一般这么定义原因是它决定了后面 csv 的列参数含义取值范围与注意点--m --n --k三个矩阵维度从 128 扫到 8192步长取 2 的幂--dtypefp32 / tf32 / fp16决定走sgemm、改 math mode 还是hgemm--iters正式迭代次数按 2.2 的规则反推不设默认值拍脑袋--no-flush关闭缓存清理只在测大矩阵或对照实验时用--csv追加一行结果必须包含形状、精度、耗时、算力四列第一次跑通的目标不是拿数字而是确认三件事程序能正常退出、CB()宏没有触发、耗时量级和手算的 FLOPs 对得上。如果 4096 的三次方跑出来只有几 GFLOPS先别怀疑硬件去看是不是把beta传成了未初始化的值或者lda传成了k而不是m。3. GEMM 基准的账要对得上column-major、leading dimension 与 TFLOPS 规则GEMM 是这套基准里最容易“数字好看但没意义”的部分。原因有三个cuBLAS 只认列主序而绝大多数人手里的矩阵是行主序leading dimension 传小了会越界读程序照样跑得飞快因为它在读别人的内存算力的分子分母如果口径不统一和别人的数据根本没法比。3.1 行主序的矩阵怎么交给列主序的 cuBLAScuBLAS 的约定是列主序内存里连续的那一维是行下标。行主序的A(m×k)在列主序视角下等价于A_cm(k×m) A_rm^T。想算C_rm A_rm * B_rm把它整体转置一次C_rm^T B_rm^T * A_rm^T左右两边换位。落到调用上就是把 M 和 N 互换、A 和 B 的位置互换// 内存按行主序存放A 是 m×kB 是 k×nC 是 m×n // 目标C_rm alpha * A_rm * B_rm beta * C_rm // 列主序视角下 C_cm B_cm * A_cm所以 m/n 互换、A/B 换位 const float alpha 1.0f, beta 0.0f; CB(cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, N, M, K, // 注意顺序是 N, M, K alpha, dB, N, // 第一个操作数现在是 Blda n dA, K, // 第二个操作数是 Aldb k beta, dC, N)); // ldc n行主序下 C 的行宽逻辑说明这个变换不改变任何数值结果只是让“内存里的行主序布局”和“cuBLAS 眼里的列主序布局”对齐省掉了显式转置的拷贝开销。对 4096 规模的方阵一次显式转置的拷贝代价足以让测出来的性能掉三成以上属于典型的自己给自己制造瓶颈。参数说明lda是第一个操作数在主维度上的行数。当transa CUBLAS_OP_N时A按列主序是M × K因此lda max(1, M)当transa CUBLAS_OP_T时lda max(1, K)。上面的调用里 A、B 换位后各自的ld就是K和N。这三个ld只要有一个传小cuBLAS 不会报错它会老老实实按你给的步长去读读到相邻矩阵的数据结果错得毫无征兆性能却正常。3.2 TFLOPS 的分子分母2MNK、缓存清理与迭代次数一次M×N×K的 GEMM 要做M*N次点积每次点积K次乘加每次乘加算两个浮点操作double flops 2.0 * (double)M * N * K; double tflops flops / (ms * 1e-3) / 1e12; // ms 是 timeit 返回的均值 // 对照用的等效带宽口径写清楚别和 GEMV 的混用 double bytes 4.0 * ((double)M * K (double)K * N (double)M * N); double gbps bytes / (ms * 1e-3) / 1e9;分子上的2是固定的M*N*K也是固定的真正会因为选择不同而产生分歧的是分母算进去的是不是只有 kernel 时间、包不包含alpha/beta的额外乘法、有没有把C的读写算进去。我一般的做法是明确声明“只统计单次 cuBLAS 调用在设备侧的耗时迭代取均值剔除缓存清理成本”并在 csv 里用一列记下flushon/off。L2 清理到底影响多大取决于矩阵工作集和 L2 容量的比值。A、B、C三块加起来小于 L2 时第二轮迭代开始数据就在片上测出来的“算力”会显著高于真实场景。这类数字拿去汇报会被质疑而且质疑得对。判据很简单把flush打开再跑一遍两次差得离谱说明之前的数字就是缓存红利。注意打开flush后小矩阵的数字会掉得很明显这不是退步而是回到了“数据必须从显存重新搬一次”的真实情况。3.3 三种精度入口SGEMM、HGEMM 与 TF32 快路同一套扫描脚本要能跑三种精度接口层必须分开写因为alpha/beta的类型跟着计算精度走。// 1) FP32 基线 const float a_f 1.0f, b_f 0.0f; CB(cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, N, M, K, a_f, dB, N, dA, K, b_f, dC, N)); // 2) FP32 输入走 TF32 快速路径先切换 math mode仍然调 sgemm CB(cublasSetMathMode(handle, CUBLAS_TF32_TENSOR_OP_MATH)); CB(cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, N, M, K, a_f, dB, N, dA, K, b_f, dC, N)); // 3) FP16alpha/beta 必须是 half 类型 __half a_h __float2half(1.0f), b_h __float2half(0.0f); CB(cublasHgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, N, M, K, a_h, dB16, N, dA16, K, b_h, dC16, N)); // 4) 显式指定计算精度cublasGemmEx便于做 A/B 对照 CB(cublasGemmEx(handle, CUBLAS_OP_N, CUBLAS_OP_N, N, M, K, a_f, dB, CUDA_R_32F, N, dA, CUDA_R_32F, K, b_f, dC, CUDA_R_32F, N, CUBLAS_COMPUTE_32F_FAST_TF32, CUBLAS_GEMM_DEFAULT));逻辑说明第 2 种和第 4 种在数学上都是“FP32 输入、低精度累加或低精度乘法”区别在于第 2 种是会话级开关一次设置对后续所有sgemm生效容易污染其他实验第 4 种是调用级作用域干净做对比实验时我更倾向用它。第 3 种的坑在alpha/beta的类型传成float指针会读出错误的位模式。参数说明CUBLAS_COMPUTE_32F是纯 FP32 累加CUBLAS_COMPUTE_32F_FAST_TF32才允许走低精度乘法单元CUBLAS_GEMM_DEFAULT是让库自己选内核不要在这里手填算法号除非你在做算法级对比。三种精度的对照关系建议固定成一张表长期放在结果目录里接口输入类型累加精度典型用途cublasSgemmFP32FP32正确性基线与数值对照cublasSgemm TF32 math modeFP32TF32训练/推理吞吐上限探测cublasHgemmFP16FP16半精度推理链路的基线cublasGemmEx可混合可选精度与吞吐的逐项对照同一组 4096 方阵FP32 基线和 TF32 路径的差距通常在数倍量级。如果跑出来只差几个百分点先别高兴去查 math mode 有没有真的生效——这是第 5 章第 3 条要展开的事。4. GEMV 是带宽题不是算力题算术强度、batched 调用与验收线GEMM 的优化目标是逼近算力峰值GEMV 的目标完全不同。把 GEMV 的 TFLOPS 拿去和 GEMM 比结论一定是“慢得离谱”但这个结论没有信息量。GEMV 的算术强度低到几乎只受显存带宽约束衡量它的指标应该是 GB/s而不是 TFLOPS。4.1 为什么 GEMV 永远无法达到峰值算力算术强度定义为浮点操作数除以搬运字节数。对y A * xA是m×n一次计算要做2mn次浮点操作至少要搬4mn字节的矩阵FP32于是算术强度 ≈ 2mn / (4 * (mn m n)) ≈ 0.5 FLOP/Byte这个值几乎不随矩阵规模变化因为m、n变大时分子分母同比放大。而硬件的“拐点”等于峰值算力除以峰值带宽现代 GPU 上这个比值通常在几十 FLOP/Byte 的量级。0.5 和几十之间差了两个数量级意味着 GEMV 无论怎么调都不可能落在算力受限区。形状m×n搬运字节FP32约浮点操作数算术强度预期瓶颈1024×10244 MB2 MFLOP~0.5带宽且可能整块进 L28192×8192256 MB134 MFLOP~0.5纯带宽4096×641 MB0.5 MFLOP~0.5带宽 kernel 启动开销所以 GEMV 的基准只有一件事值得做把实测带宽和理论带宽的比值压到尽可能高。理论带宽可以用 deviceQuery 给出的显存位宽乘以等效频率再除以 8 估算不需要额外工具。4.2 cublasSgemv 的调用与 batched 版本写法单条 GEMV 的接口很直接坑集中在lda和incx/incy上// y alpha * A * x beta * y // A 为 m×n 列主序OP_N 时 lda max(1, m) const float alpha 1.0f, beta 0.0f; CB(cublasSgemv(handle, CUBLAS_OP_N, M, N, alpha, dA, M, // lda m dx, 1, // incx beta, dy, 1)); // incy逻辑说明cublasSgemv在内部会按m的大小走不同的 kernelm较小、n较大时倾向于按行切分反之按列切分。这个选择你控制不了也不需要控制但要知道它存在——同一段代码在m64, n65536和m65536, n64两种形状下性能曲线完全不同扫描时两个方向都要跑。参数说明incx、incy是向量元素的步长单位是元素个数不是字节数传 0 会被拒绝传大于 1 的值会走非合并访问路径性能会明显下降除非你确实在做跨步取数。当一个模型里同时有几百条 GEMV 时逐条调用会被启动开销主导。两种替代写法实测差异不小// 写法一把 n1 的 GEMM 用 strided batched 一次发出去 CB(cublasSgemmStridedBatched(handle, CUBLAS_OP_N, CUBLAS_OP_N, M, 1, N, // n 固定为 1 alpha, dA, M, (long long)M * N, // 每个 batch 的 A 起点间隔 dB, N, (long long)N, // 每个 batch 的 x长度 n beta, dC, M, (long long)M, // 每个 batch 的 y长度 m batch_count)); // 写法二指针数组版适合每个 batch 的数据不在连续显存里 CB(cublasSgemvBatched(handle, CUBLAS_OP_N, M, N, alpha, dA_array, M, dx_array, 1, beta, dy_array, 1, batch_count));逻辑说明SgemmStridedBatched要求每个 batch 的数据在显存里等距排列用一个stride就能定位这种布局在推理框架里最常见SgemvBatched接受设备端的指针数组灵活但多一次指针解引用。前者对连续布局更快后者对离散布局是唯一选择。参数说明strideA传M * N是按“每个 batch 一个完整矩阵”算的如果各 batch 共享同一个矩阵多向量乘同一权重strideA要传 0这一点很容易漏漏了就会读到越界地址。4.3 用实测带宽占理论带宽的比例当验收线GEMV 的结果表里我固定只留三列形状、耗时、实测带宽外加一列“占理论带宽比例”。比例低于 60% 就说明还有优化空间通常是向量化没打开或者 batch 太小比例超过 100% 则说明数据被缓存兜住了这个数字不能对外用。批量扫描用 shell 就够不必上框架for n in 1024 2048 4096 8192 16384; do ./build/bench_gemv --m 4096 --n $n --iters 200 --csv results/gemv.csv done这里--iters给到 200 是因为 GEMV 单次耗时短迭代少的话事件分辨率不够。扫描完成后用一段脚本把 csv 画成“带宽 vs 规模”的曲线形状通常是一条先升后平的线上升段是启动开销被摊薄平台段就是真实带宽上限。5. 排查清单这类基准最容易出问题的 5 个地方跑了几十轮之后我把遇到的问题归成两类一类是计时的问题数字不可信但代码不报错一类是参数的问题数字好看但结果不对。两类都不会给你任何提示只能靠固定流程去撞。5.1 计时与同步相关的三个问题现象一小矩阵跑出来的算力超过硬件理论峰值。原因几乎总是缓存。工作集小于 L2 时第二轮迭代开始A、B、C全在片上搬运成本被抹掉测出来的就不再是显存路径。解决方式是按l2CacheSize的两倍分配 flush 缓冲每次迭代前写一遍并把 flush 成本单独测出来扣掉如果懒得做至少要在结果里标注矩阵规模和“flushoff”让别人知道这个数字的边界。现象二第一次迭代特别慢后面突然就稳了。两个原因叠加cuBLAS 句柄在首次调用时需要做内部初始化GPU 时钟也需要持续的负载才能进入高频率状态。解决方式是把 warmup 次数显式提上来我一般给 10 次warmup 结果全部丢弃并且在正式计时前插一次cudaDeviceSynchronize()确保没有残留任务干扰起点。现象三同一份二进制连跑两次结果差 15% 以上。常见原因是设备被其他进程共享、机箱散热导致降频、或者迭代总时长太短被抖动主导。解决方式是把单次迭代的总时长拉到 200 ms 以上正式结果取三轮的中位数而不是单轮均值并记录运行时的时钟频率如果条件允许独占设备跑别一边编译一边测。5.2 参数与规则相关的两个问题现象四FP32 和 TF32 只差几个百分点明显不合理。原因是低精度路径没真的生效。用cublasSgemm时忘记cublasSetMathMode(handle, CUBLAS_TF32_TENSOR_OP_MATH)或者用了cublasGemmEx但 compute type 填了CUBLAS_COMPUTE_32F都会静默回落到纯 FP32。解决方式是把 compute type 打印到日志里并固定一个对照实验同一个形状分别跑COMPUTE_32F和COMPUTE_32F_FAST_TF32两个数字应该明显分层分不开就是没生效。现象五性能正常但数值不对CPU 参考实现一比就错。最常见的是lda/ldb/ldc传小或者行主序与列主序搞混。越界读不会触发任何错误程序反而因为读到了更“近”的内存而跑得更快。解决方式是在扫描之前先做小尺寸校验取m7, n5, k3这种不规整的形状用三重循环的 CPU 版本算出参考值逐元素比对相对误差误差在 1e-5 以内才允许进入大规模扫描。这个校验代码只要二十行但能拦住后面几天的无效数据。注意这五个问题里前三个影响数字可信度后两个影响结论正确性。前三个靠流程解决后两个只能靠交叉校验没有捷径。6. 让数字站得住交叉验证、cublasLt 试探与结果表的写法一套基准跑完最尴尬的局面是有人问“你这个数字准吗”而你只能回答“我觉得准”。要避免这个局面我会在出报告前做两件事。第一件是内部交叉验证。同一个矩阵乘用cublasSgemm纯 FP32 路径和cublasGemmEx且 compute type 设为CUBLAS_COMPUTE_32F各跑一遍两者在数学上等价耗时差异应该落在 5% 以内如果差出一截说明其中一条路径选了不同的内核这个差异本身值得记录。数值上再用 CPU 参考实现对小尺寸逐元素比对避免“跑得快但算错了”这种最难查的情况。第二件是 adjudicate 一下自己写 kernel 的必要性。cuBLAS 默认给的是通用启发式选择cuBLAS 之上还有一层cublasLt接口可以拿到候选算法列表和它们各自需要的 workspace 大小逐个试一遍再取最优cublasLtMatmulHeuristicResult_t heur; int returned 0; CB(cublasLtMatmulAlgoGetHeuristic(lt_handle, lt_op, layout_a, layout_b, layout_c, layout_d, pref, heur, 1, returned)); // returned 0 时 heur.algo 才是有效候选别跳过这个判断这段代码的价值不在能提速多少而在于它能告诉你“通用启发式离最优解还有多远”。如果差距只有 3%自研 kernel 的投入基本不值得如果差 30% 以上说明这个形状确实没被库覆盖好那才是值得动手的信号。翻车点很明确returned为 0 时heur是未初始化的直接拿去调cublasLtMatmul会返回参数错误一定要先判。最后是结果表的列设计我固定成下面这样多一列都嫌累赘精度形状 (M×N×K)耗时 (ms)TFLOPS占估算峰值flushfp324096³待填待填待填ontf324096³待填待填待填onfp164096³待填待填待填onfp328192×8192×64待填待填待填off“占估算峰值”这一列是给自己看的峰值用 SM 数乘以每 SM 每时钟的 FMA 数再乘频率估算就行不必追求精确——它的作用是让你一眼看出离理论上限还差多远而不是拿去当宣传数字。我自己踩得最深的一次是拿一份没开 flush 的小矩阵数据去论证某条推理路径“已经接近硬件极限”后来被人用同样代码加上缓存清理重跑数字掉了一半多。从那以后我的习惯是任何小于 L2 容量的形状结果表里必须带flushon否则数字一律不写进结论。基准代码谁都能写难的是让报出来的每个数都能被另一个人用另一台机器复现希望帮到你。本文还有配套的精品资源点击获取