简介cublas_benchmarks 是一套面向 CUDA 高性能计算开发者与深度学习底层优化人员的基准测试代码用于评估 CuBLAS 库在 GEMM通用矩阵乘法与 GEMV矩阵向量乘法上的实际性能表现。资源包共 10 个文件以 4 个 .cu 与 4 个 .h 头文件为核心分别实现矩阵/向量数据结构、CuBLAS 函数封装与计时逻辑另附 Makefile 与 README 说明编译运行方式整体仅 13KB轻量易读。代码通过设置不同 M、N、K 规模调用对应 CuBLAS 接口并同步 GPU 后记录耗时帮助读者理解如何测量与对比不同硬件或 CUDA 版本下的算力差异。目前已有 1095 人学习下载适合需要优化矩阵运算策略、排查性能瓶颈的开发者参考借鉴。1. 从一次 GEMM 调优翻车说起这套基准到底测什么很多做 CUDA 加速的同行都有过类似经历模型推理延迟卡在某个环节直觉告诉你是矩阵乘没跑满但换了几种 tile 配置、调了几轮 block size性能曲线就是不动。问题往往不在 kernel 本身而在于你根本不知道当前硬件在 GEMM 和 GEMV 这两类操作上的真实上限在哪。cublas_benchmarks 就是干这件事的——它把 cuBLAS 里最常用的矩阵乘GEMM和矩阵向量乘GEMV拉出来用不同维度、不同数据类型、不同转置组合跑一遍给你一张能对照的性能地图。它不训练模型、不做推理纯粹是拿 cuBLAS 当被测对象测的是库函数在你机器上的实际吞吐和延迟。适合谁正在做算子优化、想确认自己手写 kernel 有没有超过 cuBLAS 的工程师以及需要给项目选型提供数据支撑的 C 开发者。你不需要懂汇编但得能编译 CUDA 程序、看得懂 GFLOPS 和带宽数字。2. 编译与运行把基准跑起来的完整链路2.1 环境依赖与编译前检查这套基准的构建方式很直接常见做法是用 CMake 或裸 Makefile 管理。我一般先确认三件事CUDA Toolkit 版本、GPU 计算能力compute capability、以及 cuBLAS 是否随 Toolkit 一起装好了。cuBLAS 从 CUDA 10 开始就是 Toolkit 的标配组件不需要单独下载但头文件cublas_v2.h和库文件libcublas.so的路径得让编译器找得到。如果你用的是容器环境注意 nvidia-container-toolkit 要装对否则nvidia-smi能看到卡但 CUDA 程序跑不起来。编译前先跑一条命令确认环境nvcc --version nvidia-smi --query-gpuname,compute_cap,memory.total --formatcsv这条命令输出三样东西CUDA 编译器版本、GPU 型号和计算能力、显存大小。计算能力决定了你能用哪些 cuBLAS 特性比如 TF32 需要 compute capability 8.0 以上Ampere 架构起。如果nvcc找不到检查PATH里有没有/usr/local/cuda/bin如果nvidia-smi报错先解决驱动问题别急着编译。2.2 编译命令与关键编译选项假设源码目录结构是src/放.cu文件、CMakeLists.txt在根目录典型的编译流程如下mkdir -p build cd build cmake .. -DCMAKE_CUDA_ARCHITECTURES80 -DCMAKE_BUILD_TYPERelease make -j$(nproc)这里有几个参数值得展开说。CMAKE_CUDA_ARCHITECTURES80指定目标架构为 sm_80对应 Ampere。如果你用的是其他架构改成对应数字Turing 是 75Ada Lovelace 是 89Hopper 是 90。不指定的话 CMake 可能会编译一堆架构编译时间翻倍但性能没提升。CMAKE_BUILD_TYPERelease会开-O3优化基准测试必须用 ReleaseDebug 模式下 cuBLAS 调用本身不受影响但你的计时逻辑和内存分配可能被拖慢。如果项目用的是裸 Makefile核心编译命令类似nvcc -O3 -archsm_80 -lcublas -lcudart -o gemm_bench src/gemm_bench.cu-lcublas和-lcudart是必须链接的库。-archsm_80和 CMake 里的架构参数等价。注意-O3对 host 代码的优化device 代码的优化由-arch和--use_fast_math等选项控制但基准测试不建议开--use_fast_math因为它会改变浮点运算语义测出来的数字和实际生产环境对不上。2.3 运行参数与输出解读编译产物通常是一个可执行文件运行时通过命令行参数控制测试规模。常见参数包括矩阵维度M/N/K、数据类型s单精度、d双精度、c单精度复数、z双精度复数、以及操作类型gemm或gemv。一个典型的运行命令./gemm_bench --m 4096 --n 4096 --k 4096 --dtype s --op gemm --iters 100输出一般包含三列维度、耗时毫秒、计算吞吐GFLOPS。GFLOPS 的算法是2*M*N*K / (耗时 * 1e9)GEMV 则是2*M*N / (耗时 * 1e9)。看结果时注意两点第一小维度下 GFLOPS 偏低是正常的因为 kernel launch 开销和内存访问延迟占比高第二GEMV 是内存带宽瓶颈型操作GFLOPS 数字不会太高但有效带宽GB/s才是关键指标。如果你看到某个维度下 GFLOPS 突然掉一半大概率是矩阵尺寸跨过了某个 cache 层级边界比如从 L2 驻留变成 HBM 反复读写。提示跑基准前先把 GPU 时钟锁定到固定频率否则动态调频会让不同轮次的数字没法对比。用nvidia-smi -lgc可以锁核心频率具体频率值查官方规格。3. 参数怎么设维度、数据类型与转置组合的实测差异3.1 矩阵维度对 GEMM 性能的影响规律GEMM 的性能对维度非常敏感但不是线性的。我实测下来规律大致分三段。第一段是 M/N/K 都小于 512 时GFLOPS 随维度增长很快因为并行度不够SM 占用率低。第二段是 512 到 4096 之间性能爬升变缓逐渐逼近硬件的理论峰值。第三段是超过 4096 之后性能基本持平偶尔因为显存带宽争抢出现小幅波动。这里有个容易踩的坑很多人只测方阵MNK但实际业务里矩阵往往是扁的。比如全连接层里 M 是 batch size、N 是输出维度、K 是输入维度三者比例可能差一个数量级。cuBLAS 对不同形状的矩阵会走不同的内部 kernel方阵跑得好不代表扁矩阵也跑得好。建议至少测三组形状方阵4096×4096×4096、瘦长矩阵M1, N4096, K4096模拟 GEMV 场景、以及宽扁矩阵M4096, N64, K4096。# 方阵 ./gemm_bench --m 4096 --n 4096 --k 4096 --dtype s # 瘦长矩阵接近 GEMV ./gemm_bench --m 1 --n 4096 --k 4096 --dtype s # 宽扁矩阵 ./gemm_bench --m 4096 --n 64 --k 4096 --dtype s这三组跑完你对自己硬件在不同形状下的表现就有底了。如果瘦长矩阵的 GFLOPS 只有方阵的十分之一别慌这是正常的——M1 时并行度天然受限cuBLAS 会退化成类似 GEMV 的实现。3.2 数据类型与计算精度的取舍cuBLAS 支持的数据类型直接决定性能上限。单精度FP32是基准测试的默认选项但如果你在跑深度学习负载TF32 和 FP16 才是更常见的格式。TF32 是 Ampere 引入的格式在 cuBLAS 里通过CUBLAS_TF32_TENSOR_OP_MATH模式启用它用 19 位表示尾数FP32 是 23 位精度略降但吞吐能到 FP32 的数倍。FP16 则更进一步配合 Tensor Core 能跑到理论峰值的很大比例。测不同类型时注意基准程序是否支持--dtype参数切换。如果不支持你可能需要改源码里的cublasSetMathMode调用。常见做法是在初始化 handle 之后加一行cublasSetMathMode(handle, CUBLAS_TF32_TENSOR_OP_MATH);这行代码开启 TF32 模式之后所有 FP32 的 GEMM 调用都会尝试走 Tensor Core。但要注意TF32 的结果和纯 FP32 有微小差异如果你的业务对精度敏感比如科学计算别开这个模式。基准测试里可以两种模式都跑一遍对比 GFLOPS 和数值误差。3.3 转置组合与内存布局的隐藏成本cuBLAS 的 GEMM 接口支持四种转置组合NN、NT、TN、TT第一个字母是 A 的转置状态第二个是 B 的。不同组合的性能差异可能达到 20% 以上原因是 cuBLAS 内部对每种组合走了不同的 kernel 实现内存访问模式也不同。NN 是最常见的组合A 和 B 都是行主序存储TN 在深度学习里也常见因为权重矩阵通常需要转置。测转置组合时基准程序一般通过--transa和--transb参数控制。如果没有这些参数你得手动改源码里的CUBLAS_OP_N/CUBLAS_OP_T枚举值。我建议至少跑 NN 和 TN 两组因为这两组覆盖了大多数实际场景。如果 TN 比 NN 慢很多检查你的矩阵存储顺序是否和 cuBLAS 的预期一致——cuBLAS 默认列主序如果你传的是行主序矩阵实际上相当于做了一次转置性能特征会变。注意转置操作本身不复制数据只是改变索引方式所以它对性能的影响来自内存访问的连续性而不是额外的拷贝开销。4. 避坑与排查基准测试里最容易翻车的五个点4.1 计时方式不对测出来全是噪声现象同一组参数跑十次GFLOPS 波动超过 30%。原因用了std::chrono在 host 端计时但 CUDA kernel 是异步执行的host 端计时器在 kernel 还没跑完就停了。解决用 CUDA Event 计时在 kernel 前后各插一个 event然后cudaEventSynchronize等待完成再取时间差。代码模板如下cudaEvent_t start, stop; cudaEventCreate(start); cudaEventCreate(stop); cudaEventRecord(start); // 调用 cuBLAS GEMM cublasSgemm(handle, ...); cudaEventRecord(stop); cudaEventSynchronize(stop); float ms 0; cudaEventElapsedTime(ms, start, stop);cudaEventElapsedTime返回的是毫秒精度足够。注意 event 要复用别在循环里反复 create/destroy那本身就有开销。4.2 显存分配拖慢首次调用现象第一次跑某个维度特别慢后面就正常了。原因cuBLAS 在首次调用时会做内部初始化包括加载 kernel 镜像、分配 workspace。解决在正式计时前先跑几轮 warm-up通常 3 到 5 轮就够了。基准程序如果有--warmup参数就直接用没有的话在计时循环前手动加几轮空跑。4.3 矩阵尺寸没对齐性能断崖式下跌现象M4095 和 M4096 的 GFLOPS 差一倍。原因cuBLAS 内部对矩阵维度有对齐要求通常是 8 或 16 的倍数。不对齐时kernel 会走边界处理分支效率大幅下降。解决测基准时尽量用 2 的幂次或 64 的倍数作为维度。如果业务场景确实有不规则维度单独测一组并记录别和规则维度混在一起看平均值。4.4 多卡环境没指定设备跑错卡现象结果和预期硬件规格对不上比如 A100 跑出 V100 的数字。原因程序默认用 device 0但机器上 device 0 可能不是你预期的那张卡。解决在代码开头加cudaSetDevice(device_id)或者用CUDA_VISIBLE_DEVICES环境变量控制可见设备。跑之前先用nvidia-smi -L确认设备编号和型号的对应关系。4.5 忽略功耗墙和温度墙现象跑了几分钟之后 GFLOPS 逐渐下降。原因GPU 撞到功耗墙或温度墙自动降频。解决跑长时间基准前先监控nvidia-smi -q -d POWER,TEMPERATURE确认功耗和温度在安全范围内。如果散热条件有限缩短单次测试时间或者分批跑、中间留冷却间隔。数据中心卡如 A100的功耗墙通常可以通过nvidia-smi -pl调整但消费级卡往往锁死只能改善散热。5. 进阶技巧用基准数据反推 kernel 优化方向跑完基准拿到一堆 GFLOPS 数字之后真正的价值在于对比。我一般会做三件事。第一把 cuBLAS 的结果和自己手写 kernel 的结果放在同一张表里看差距在哪。如果手写 kernel 在某个维度下只有 cuBLAS 的 60%先别急着改 tile 大小用 Nsight Compute 抓一下 memory throughput 和 compute throughput看瓶颈在访存还是计算。第二把不同数据类型的数字横向对比如果 FP16 只有 FP32 的 1.5 倍而不是预期的 4 倍以上说明 Tensor Core 没吃满检查矩阵维度是否满足 Tensor Core 的对齐要求通常是 16 的倍数。第三把 GEMV 的有效带宽算出来和显存理论带宽对比如果不到 70%说明访存模式有问题考虑合并访问或使用向量化加载。下面这张表是我常用的对照模板你可以照着填自己的数据操作类型维度 (M×N×K)数据类型耗时 (ms)GFLOPS有效带宽 (GB/s)GEMM4096³FP32待填待填不适用GEMM4096³TF32待填待填不适用GEMM4096³FP16待填待填不适用GEMV4096×4096FP32待填待填待填填完这张表你对自己硬件的脾气就摸得差不多了。从那以后我每次做算子优化前都强制走一遍基准流程先锁频、再 warm-up、然后按形状和类型矩阵式跑一遍最后拿数据说话。这套习惯帮我省下了大量“凭感觉调参”的时间也避免了很多次“以为优化了其实只是噪声”的翻车。希望帮到你。本文还有配套的精品资源点击获取