CUDA矩阵乘性能基准测试：cuBLAS GEMM/GEMV调优与避坑指南-育师

简介：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-gpu=name,compute_cap,memory.total --format=csv

这条命令输出三样东西：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_ARCHITECTURES=80 -DCMAKE_BUILD_TYPE=Release make -j$(nproc)

这里有几个参数值得展开说。CMAKE_CUDA_ARCHITECTURES=80指定目标架构为 sm_80，对应 Ampere。如果你用的是其他架构，改成对应数字：Turing 是 75，Ada Lovelace 是 89，Hopper 是 90。不指定的话 CMake 可能会编译一堆架构，编译时间翻倍但性能没提升。CMAKE_BUILD_TYPE=Release会开-O3优化，基准测试必须用 Release，Debug 模式下 cuBLAS 调用本身不受影响，但你的计时逻辑和内存分配可能被拖慢。

如果项目用的是裸 Makefile，核心编译命令类似：

nvcc -O3 -arch=sm_80 -lcublas -lcudart -o gemm_bench src/gemm_bench.cu

-lcublas和-lcudart是必须链接的库。-arch=sm_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 之后，性能基本持平，偶尔因为显存带宽争抢出现小幅波动。

这里有个容易踩的坑：很多人只测方阵（M=N=K），但实际业务里矩阵往往是扁的。比如全连接层里 M 是 batch size、N 是输出维度、K 是输入维度，三者比例可能差一个数量级。cuBLAS 对不同形状的矩阵会走不同的内部 kernel，方阵跑得好不代表扁矩阵也跑得好。建议至少测三组形状：方阵（4096×4096×4096）、瘦长矩阵（M=1, N=4096, K=4096，模拟 GEMV 场景）、以及宽扁矩阵（M=4096, N=64, K=4096）。

# 方阵 ./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 只有方阵的十分之一，别慌，这是正常的——M=1 时并行度天然受限，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 矩阵尺寸没对齐，性能断崖式下跌

现象：M=4095 和 M=4096 的 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%，说明访存模式有问题，考虑合并访问或使用向量化加载。

下面这张表是我常用的对照模板，你可以照着填自己的数据：