简介CuBLAS是NVIDIA提供的GPU线性代数库其中的GEMM与GEMV分别指通用矩阵乘法和矩阵向量乘法是深度学习、科学计算中频繁出现的高频核心算子。这份cublas_benchmarks为CUDA开发者准备了一套可直接编译运行的基准测试程序重点面向需要评估CuBLAS性能、对比不同硬件配置或进行矩阵运算调优的工程师与研究人员。压缩包内共10个文件大小约13KB包含4个C头文件、4个CUDA源文件以及构建脚本和说明文档头文件负责统一声明矩阵工具与CuBLAS函数接口CUDA源文件分别实现数据初始化、GEMM/GEMV调用与GPU同步计时构建脚本则简化了编译流程。已有1095人学习下载。通过调整矩阵尺寸参数并记录执行时间读者可以直观了解不同运算规模下CuBLAS的实际吞吐量还能自行扩展测试维度为后续代码优化和计算策略选择提供可靠的实测依据。1. 性能基准不是跑个分就完事cublas_benchmarks 到底在测什么做 GPU 矩阵计算优化的人最烦的就是别人问“你到底快了多少”。快没快不能靠感觉也不能只盯着某个 kernel 的局部时间。cuBLAS 里的 GEMM 和 GEMV 几乎是一切线性代数应用的地基矩阵分解、神经网络全连接层、迭代求解器全在它们上面跑。cublas_benchmarks 这套基准代码就是把不同尺寸、不同转置组合、不同款数据类型的矩阵乘法和矩阵向量乘法挨个跑一遍输出可对比的 TFLOPS 和有效带宽让你在动手写自己的 kernel 之前先知道硬件真正的天花板在哪。适合做 GPU 计算、性能调优、以及在 C 项目里集成 cuBLAS 的开发者。新手拿它当标准测试工具熟手改参数测自己业务场景都很顺手。2. GEMM 与 GEMV 的性能模型为什么基准要先算一遍理论峰值2.1 从计算量和访存量推导 FLOPs 与带宽基准测试不是简单丢一个矩阵进去跑出个数字。你得先知道理论上限在哪否则结果好坏无从判断。GEMM 的公式是 C alpha * A * B beta * C其中 A 是 M x KB 是 K x NC 是 M x N。浮点计算量大约为 2 * M * N * K这个“2”来自乘法和加法各一次。GEMV 则是 y alpha * A * x beta * y计算量是 2 * M * K比 GEMM 低了一个量级。别看计算量低GEMV 的访存特征完全不同——它要读整个 A 矩阵但每个元素只用一次所以性能瓶颈几乎都在显存带宽。理论峰值怎么算比如某块 GPU 的 FP32 计算峰值是 15 TFLOPS那 GEMM 的实际 TFLOPS 达到 80% 就算正常。GEMV 则要对比 HBM 带宽比如显存带宽是 900 GB/s你的 GEMV 有效带宽能到 700 GB/s 就不错了。基准代码里通常会把这两个指标都打出来而且会附上计算强度FLOPs / Bytes方便你判断当前测试是计算密集还是访存密集。2.2 影响 cuBLAS 性能的四个关键参数第一个是矩阵尺寸 M、N、K。K 维尤其敏感因为它决定累加次数K 较小时 compute-to-memory 比率会掉甚至让 GEMM 退化成访存密集型。第二个是转置设置。cuBLAS 默认列主序CUBLAS_OP_T 和 CUBLAS_OP_N 的组合不同内核里访问 A 和 B 的模式就不同。同一个尺寸转置选错可能直接掉 30% 性能。第三个是数据类型。FP16 用上 Tensor Core 能达到 FP32 的好几倍而 FP64 可能只有 FP32 的一半。第四个是内存对齐方式。用 cudaMalloc 分配的显存对齐到 256 字节基准测试才能吃到最优路径如果传入的是某些托管内存或者非对齐的偏移地址cuBLAS 会走通用 kernel性能断崖式下降。2.3 基准脚本的测试矩阵设计常见做法是让矩阵尺寸按对数分布取比如从 64 到 8192每次乘以 1.5 或者按 2 的幂次跳。为什么要这样因为 GPU 内核对尺寸的敏感度不是线性的某些尺寸正好落在 tile 边界上性能极高偏一个数就可能掉下来。跑的时候对每个配置做多次迭代取最短时间而不是平均时间——短时间才代表内核真正跑满状态平均时间会被偶发降频污染。2.4 从基准结果反推 cuBLAS 内部策略当你跑完 cublas_benchmarks 之后会得到一张表M、N、K、TFLOPS、带宽。这张表能反推出很多信息。比如当 M 和 N 很小但 K 很大时TFLOPS 通常很低原因是矩阵太小GPU 的并行度喂不饱。这时 cuBLAS 可能会切换到另一种 kernel 策略展开循环减少同步开销。反过来当 M 和 N 都很大时TFLOPS 会逼近理论峰值说明 tile 切分和寄存器复用做得很好。看表的时候重点关注“性能拐点”的位置那往往是你的算法需要避开的尺寸区间。3. cublas_benchmarks 源码拆解从调用到计时再到输出3.1 项目文件结构与主流程这套基准一般会包含main.cpp、bench_gemm.cu、bench_gemv.cu、common.h这几个文件。common.h里定义了测试配置结构体比如尺寸、转置、数据类型、迭代次数。bench_gemm.cu专门封装 cublasGemm 的调用和计时逻辑bench_gemv.cu则是对应的 GEMV 版本。主流程在main.cpp里做的事很简单解析命令行参数根据参数分配显存初始化随机数据然后进入每个测试配置的循环。struct BenchConfig { int m, n, k; cublasOperation_t transA, transB; cudaDataType_t dataType; int iterations; int warmup; };这段结构体是整套基准的核心。m/n/k不必多说transA/transB控制是否转置dataType支持CUDA_R_32F、CUDA_R_16F等类型。iterations是计时循环次数warmup是预热次数。我把这些字段全放在一个结构体里就是为了后续扩展方便——比如要测批量 GEMM直接加一个batchCount字段就行。3.2 核心计时方法CUDA 事件与流同步计时的正确姿势是用cudaEventRecord而不是clock()或者std::chrono。GPU 运行是异步的CPU 计时器根本不知道内核什么时候结束。标准做法是在同一个流上记录开始事件、运行内核、记录结束事件然后cudaEventSynchronize等待结束事件完成最后取两个事件的时间差。cudaEvent_t start, stop; cudaEventCreate(start); cudaEventCreate(stop); cudaEventRecord(start, stream); for (int i 0; i config.iterations; i) { cublasSgemm(handle, config.transA, config.transB, config.m, config.n, config.k, alpha, dA, lda, dB, ldb, beta, dC, ldc); } cudaEventRecord(stop, stream); cudaEventSynchronize(stop); float ms 0.0f; cudaEventElapsedTime(ms, start, stop);注意循环内只调用 cublas 函数没有别的操作。循环结束后统一同步这样能摊薄事件本身的调用开销。时间ms是全部迭代的总耗时最后算单次耗时需要除以iterations。这个设计有个好处如果某个 kernel 启动太快单次事件分辨率不够循环多次后误差会降到很低。3.3 参数解析与可扩展性命令行解析部分我习惯用getopt或简单的字符串匹配不引第三方依赖。关键参数要覆盖--m、--n、--k、--transa、--transb、--dtype、--iterations、--warmup。如果没传某个参数就给默认值比如尺寸 1024FP32不转置。解析完参数后把配置填进BenchConfig再调用对应的 bench 函数。这种设计让新加一种数据类型变得很容易只要在switch里加一个 case。nvcc -O3 -archsm_80 -o cublas_bench bench_gemm.cu bench_gemv.cu main.cpp -lcublas编译时注意-arch要跟你实际的 GPU 计算能力匹配。sm_80是 Ampere 世代常见选择如果显卡更新需要用sm_86或sm_90。链接必须加-lcublas否则一堆未定义符号。跑的时候也简单./cublas_bench --m 2048 --n 2048 --k 2048 --dtype fp32 --iterations 100 --warmup 103.4 数据初始化与正确性校验性能测试最容易翻车的一点是为了省时间初始化矩阵时全部填 0。这会导致 cuBLAS 内部可能走快速路径比如检测到 beta0 就直接跳过写 C或者某些 kernel 针对稀疏结构做了优化测出来的数字虚高。我的做法是每个矩阵用随机数填满但不上cudaMemset后再覆盖一部分。随机数生成用curand或者 CPU 生成后拷过去。基准代码里通常会留一个--check开关打开后会把 GPU 结果拷回 CPU和朴素的 CPU 三重循环结果对比误差超过阈值就报错。// 初始化 A 和 B 为随机数C 为随机数避免 beta 路径被跳过 for (int i 0; i totalA; i) { hA[i] (float)rand() / RAND_MAX; }要注意随机数种子固定下来这样每次跑的结果可复现。有些人在每次运行时用time(NULL)做种子结果跑出来的性能波动还以为是自己优化出了问题其实只是数值变了导致里面的除法分支不同。固定种子是基准测试的规矩不是可有可无。4. 基准测试避坑指南为什么你的数字比别人低一半4.1 现象TFLOPS 只有理论峰值的 30%我第一次拿这套基准跑旧卡就是这样以为卡坏了后来发现是忘了开 Tensor Core。cuBLAS 默认的cublasSgemm在 FP32 下走的是普通 CUDA core 路径并不会自动用上 TF32 加速。要跑到更高性能得用cublasGemmEx配合CUBLAS_COMPUTE_32F_FAST_TF32。原因就是 cuBLAS 为了兼容性和精度把数学模式设成了默认的精确模式除非你显式放宽精度要求。解决方法是换用cublasGemmEx并把计算类型设为CUBLAS_COMPUTE_32F_FAST_TF32或者直接用 FP16 输入让内核自动走 Tensor Core。从那以后我跑基准之前都会先查一遍设备的计算能力确认哪种精度组合能吃满硬件。4.2 现象GEMV 带宽远低于显存带宽GEMV 的本质是访存密集型按理说带宽应该能打得很高。但很多人测出来连一半都不到。原因在于 A 矩阵是列主序存储的而 GEMV 的访问模式是逐行读取每个元素做乘加。如果trans参数设置不当相当于每次访存都跨越一整列局部性极差。解决方法是测两种转置模式CUBLAS_OP_N和CUBLAS_OP_T选带宽高的那个。另外当 M 和 K 都很小的时候GEMV 的启动开销占了大头合理做法是改用cublasSgemvBatched把多个向量一次性交给 GPU摊薄启动延迟。4.3 现象计时结果波动大同一配置两次跑差 20%这不是玄学是 GPU 的实时时钟频率在捣鬼。GPU 会根据负载动态升频和降频你跑基准时如果散热不好功耗墙就触发降频结果自然飘。解决方法是三管齐下一是开足--warmup次数让核心加热后进入稳定频率再计时二是增加迭代次数比如 200 次以上然后取最小值三是在运行前用nvidia-smi锁定时钟频率。锁时钟需要 root 权限如果没有至少要让测试程序反复运行同一个 kernel让频率稳定下来。4.4 现象换了一台 GPU 后相对性能排序完全颠倒同样的 cublas_benchmarks在 A 卡上 FP16 比 FP32 快 4 倍到 B 卡上只快 2 倍。原因是不同世代的硬件对 Tensor Core 的调度策略差异很大cuBLAS 内部维护着多个 kernel 变体根据尺寸、对齐和硬件代数自动选择。所以跨卡比较时不能只看 TFLOPS 高低还要把驱动的 CUDA 版本、cuBLAS 版本、卡的计算能力打出来否则你只能看到结果看不到背后的选路逻辑。我的习惯是把这些环境信息写进结果文件头部这样复现的时候不会一头雾水。4.5 现象第一次调用某个尺寸时特别慢后面就正常了这是 cuBLAS 的 lazy loading 机制在起作用。第一次调用某个 API 时它会加载对应 kernel 模块做各种 heuristics 搜索比如试探多个 tile 大小来选择最优配置这个开销可能高达几十毫秒。而正式的迭代循环只有几百微秒如果直接把第一次调用也计时最终平均时间会被严重拉高。解决方法是一定要做预热预热次数最好覆盖到每个不同的尺寸和转置组合。代码里可以设计成每个配置单独预热或者干脆跑一遍完整的空测试把所有 kernel 策略都触发一遍。5. 进阶让基准结果反映你的真实场景5.1 从 GEMM 扩展到 StridedBatch GEMM如果你的实际场景是小矩阵批量运算比如图神经网络里几百个邻接矩阵的乘法单测一个 GEMM 没有意义。cublas_benchmarks 这种框架很容易扩展出分块批量模式。常见做法是加一个--batch参数然后调用cublasStridedBatchedGemmEx。关键是要正确设置 stride 参数。对于连续存储的批量矩阵A 的 stride 是M*KB 的 stride 是K*NC 的 stride 是M*N。错一个 stride结果就是错的而且错得很有隐蔽性因为每个 batch 内部可能还算出了一个“合理”的数字。cublasStridedBatchedGemmEx(handle, transA, transB, m, n, k, alpha, dA, lda, strideA, dB, ldb, strideB, beta, dC, ldc, strideC, batchCount, computeType, CUBLAS_GEMM_DEFAULT_TENSOR_OP);参数里strideA、strideB、strideC都是元素个数不是字节数。这块最容易混。我吃过一次亏把字节数填进去结果前 256 个 batch 数据错了后面的因为内存越界直接报错。多看几遍文档或者用一个小 batch 尺寸验证数值能省下很多 debug 时间。批量模式下性能曲线和单发 GEMM 完全不同小尺寸时批量越大吞吐越高直到批量带来的调度开销超过并行收益。5.2 用 Roofline 模型解释基准数据拿到基准结果后我通常会先把每个点按计算强度分类画一条简单的手工 Roofline。不需要画图工具列表就行。对于 GEMM计算强度 2MNK / (MN NK MK)当 MNK4096 时强度大约为 1365 FLOPs/Byte远超机器平衡点属于计算密集。此时你做内存优化没用得靠提高计算密度比如换 Tensor Core。对于 GEMV计算强度只有 2K / (M*K ...) 约等于 2/M 当 M 很大时妥妥的访存密集。这时候锁频、调转置比什么都管用。基准代码里如果能输出计算强度整个性能分析会直观很多。5.3 自定义一个与你业务相同的测试配置标准基准测的是方阵和幂等尺寸但你的业务可能是 M1024、N128、K768 这种非对称尺寸。不用改代码只要在调用入口传入这些值即可。我见过有人为了测某个深度优化的卷积等效 GEMM直接把m/n/k设成卷积输出的形状参数这样能提前判断 cuBLAS 的通用实现是否够用还是必须自己写 kernel。如果你想测得更细可以在BenchConfig里加一个alpha和beta的字段。很多人默认 alpha1、beta0但真实场景里 beta 往往是 1也就是累加操作。cuBLAS 对 beta0 和 beta1 可能会走不同的 kernel 路径beta1 时可能要多读一次 C 再写回性能差几个百分点。5.4 验证基准本身没有测错最后一个要养成的习惯是每次新增测试场景时先跑一遍带--check的验证模式确认结果数值正确再开始谈性能。我过去有一次为了追求速度把迭代次数加到 2000结果矩阵尺寸太大导致数据溢出TFLOPS 看起来很高实际算出的 C 已经全是无穷大。你要是在这种错误数据上做优化后面所有优化决策都会跟着错。验证模式的做法很简单在 CPU 上用双精度跑一遍同样的运算和 GPU 结果比较相对误差小于 1e-3 就通过。从那以后我每次做性能优化都会强制先跑一遍基准并验证数值确认环境没问题再看性能曲线。这套流程虽然多花两分钟但能省下后面一小时的 debug 时间。希望帮到你。本文还有配套的精品资源点击获取
简介CuBLAS是NVIDIA提供的GPU线性代数库其中的GEMM与GEMV分别指通用矩阵乘法和矩阵向量乘法是深度学习、科学计算中频繁出现的高频核心算子。这份cublas_benchmarks为CUDA开发者准备了一套可直接编译运行的基准测试程序重点面向需要评估CuBLAS性能、对比不同硬件配置或进行矩阵运算调优的工程师与研究人员。压缩包内共10个文件大小约13KB包含4个C头文件、4个CUDA源文件以及构建脚本和说明文档头文件负责统一声明矩阵工具与CuBLAS函数接口CUDA源文件分别实现数据初始化、GEMM/GEMV调用与GPU同步计时构建脚本则简化了编译流程。已有1095人学习下载。通过调整矩阵尺寸参数并记录执行时间读者可以直观了解不同运算规模下CuBLAS的实际吞吐量还能自行扩展测试维度为后续代码优化和计算策略选择提供可靠的实测依据。1. 性能基准不是跑个分就完事cublas_benchmarks 到底在测什么做 GPU 矩阵计算优化的人最烦的就是别人问“你到底快了多少”。快没快不能靠感觉也不能只盯着某个 kernel 的局部时间。cuBLAS 里的 GEMM 和 GEMV 几乎是一切线性代数应用的地基矩阵分解、神经网络全连接层、迭代求解器全在它们上面跑。cublas_benchmarks 这套基准代码就是把不同尺寸、不同转置组合、不同款数据类型的矩阵乘法和矩阵向量乘法挨个跑一遍输出可对比的 TFLOPS 和有效带宽让你在动手写自己的 kernel 之前先知道硬件真正的天花板在哪。适合做 GPU 计算、性能调优、以及在 C 项目里集成 cuBLAS 的开发者。新手拿它当标准测试工具熟手改参数测自己业务场景都很顺手。2. GEMM 与 GEMV 的性能模型为什么基准要先算一遍理论峰值2.1 从计算量和访存量推导 FLOPs 与带宽基准测试不是简单丢一个矩阵进去跑出个数字。你得先知道理论上限在哪否则结果好坏无从判断。GEMM 的公式是 C alpha * A * B beta * C其中 A 是 M x KB 是 K x NC 是 M x N。浮点计算量大约为 2 * M * N * K这个“2”来自乘法和加法各一次。GEMV 则是 y alpha * A * x beta * y计算量是 2 * M * K比 GEMM 低了一个量级。别看计算量低GEMV 的访存特征完全不同——它要读整个 A 矩阵但每个元素只用一次所以性能瓶颈几乎都在显存带宽。理论峰值怎么算比如某块 GPU 的 FP32 计算峰值是 15 TFLOPS那 GEMM 的实际 TFLOPS 达到 80% 就算正常。GEMV 则要对比 HBM 带宽比如显存带宽是 900 GB/s你的 GEMV 有效带宽能到 700 GB/s 就不错了。基准代码里通常会把这两个指标都打出来而且会附上计算强度FLOPs / Bytes方便你判断当前测试是计算密集还是访存密集。2.2 影响 cuBLAS 性能的四个关键参数第一个是矩阵尺寸 M、N、K。K 维尤其敏感因为它决定累加次数K 较小时 compute-to-memory 比率会掉甚至让 GEMM 退化成访存密集型。第二个是转置设置。cuBLAS 默认列主序CUBLAS_OP_T 和 CUBLAS_OP_N 的组合不同内核里访问 A 和 B 的模式就不同。同一个尺寸转置选错可能直接掉 30% 性能。第三个是数据类型。FP16 用上 Tensor Core 能达到 FP32 的好几倍而 FP64 可能只有 FP32 的一半。第四个是内存对齐方式。用 cudaMalloc 分配的显存对齐到 256 字节基准测试才能吃到最优路径如果传入的是某些托管内存或者非对齐的偏移地址cuBLAS 会走通用 kernel性能断崖式下降。2.3 基准脚本的测试矩阵设计常见做法是让矩阵尺寸按对数分布取比如从 64 到 8192每次乘以 1.5 或者按 2 的幂次跳。为什么要这样因为 GPU 内核对尺寸的敏感度不是线性的某些尺寸正好落在 tile 边界上性能极高偏一个数就可能掉下来。跑的时候对每个配置做多次迭代取最短时间而不是平均时间——短时间才代表内核真正跑满状态平均时间会被偶发降频污染。2.4 从基准结果反推 cuBLAS 内部策略当你跑完 cublas_benchmarks 之后会得到一张表M、N、K、TFLOPS、带宽。这张表能反推出很多信息。比如当 M 和 N 很小但 K 很大时TFLOPS 通常很低原因是矩阵太小GPU 的并行度喂不饱。这时 cuBLAS 可能会切换到另一种 kernel 策略展开循环减少同步开销。反过来当 M 和 N 都很大时TFLOPS 会逼近理论峰值说明 tile 切分和寄存器复用做得很好。看表的时候重点关注“性能拐点”的位置那往往是你的算法需要避开的尺寸区间。3. cublas_benchmarks 源码拆解从调用到计时再到输出3.1 项目文件结构与主流程这套基准一般会包含main.cpp、bench_gemm.cu、bench_gemv.cu、common.h这几个文件。common.h里定义了测试配置结构体比如尺寸、转置、数据类型、迭代次数。bench_gemm.cu专门封装 cublasGemm 的调用和计时逻辑bench_gemv.cu则是对应的 GEMV 版本。主流程在main.cpp里做的事很简单解析命令行参数根据参数分配显存初始化随机数据然后进入每个测试配置的循环。struct BenchConfig { int m, n, k; cublasOperation_t transA, transB; cudaDataType_t dataType; int iterations; int warmup; };这段结构体是整套基准的核心。m/n/k不必多说transA/transB控制是否转置dataType支持CUDA_R_32F、CUDA_R_16F等类型。iterations是计时循环次数warmup是预热次数。我把这些字段全放在一个结构体里就是为了后续扩展方便——比如要测批量 GEMM直接加一个batchCount字段就行。3.2 核心计时方法CUDA 事件与流同步计时的正确姿势是用cudaEventRecord而不是clock()或者std::chrono。GPU 运行是异步的CPU 计时器根本不知道内核什么时候结束。标准做法是在同一个流上记录开始事件、运行内核、记录结束事件然后cudaEventSynchronize等待结束事件完成最后取两个事件的时间差。cudaEvent_t start, stop; cudaEventCreate(start); cudaEventCreate(stop); cudaEventRecord(start, stream); for (int i 0; i config.iterations; i) { cublasSgemm(handle, config.transA, config.transB, config.m, config.n, config.k, alpha, dA, lda, dB, ldb, beta, dC, ldc); } cudaEventRecord(stop, stream); cudaEventSynchronize(stop); float ms 0.0f; cudaEventElapsedTime(ms, start, stop);注意循环内只调用 cublas 函数没有别的操作。循环结束后统一同步这样能摊薄事件本身的调用开销。时间ms是全部迭代的总耗时最后算单次耗时需要除以iterations。这个设计有个好处如果某个 kernel 启动太快单次事件分辨率不够循环多次后误差会降到很低。3.3 参数解析与可扩展性命令行解析部分我习惯用getopt或简单的字符串匹配不引第三方依赖。关键参数要覆盖--m、--n、--k、--transa、--transb、--dtype、--iterations、--warmup。如果没传某个参数就给默认值比如尺寸 1024FP32不转置。解析完参数后把配置填进BenchConfig再调用对应的 bench 函数。这种设计让新加一种数据类型变得很容易只要在switch里加一个 case。nvcc -O3 -archsm_80 -o cublas_bench bench_gemm.cu bench_gemv.cu main.cpp -lcublas编译时注意-arch要跟你实际的 GPU 计算能力匹配。sm_80是 Ampere 世代常见选择如果显卡更新需要用sm_86或sm_90。链接必须加-lcublas否则一堆未定义符号。跑的时候也简单./cublas_bench --m 2048 --n 2048 --k 2048 --dtype fp32 --iterations 100 --warmup 103.4 数据初始化与正确性校验性能测试最容易翻车的一点是为了省时间初始化矩阵时全部填 0。这会导致 cuBLAS 内部可能走快速路径比如检测到 beta0 就直接跳过写 C或者某些 kernel 针对稀疏结构做了优化测出来的数字虚高。我的做法是每个矩阵用随机数填满但不上cudaMemset后再覆盖一部分。随机数生成用curand或者 CPU 生成后拷过去。基准代码里通常会留一个--check开关打开后会把 GPU 结果拷回 CPU和朴素的 CPU 三重循环结果对比误差超过阈值就报错。// 初始化 A 和 B 为随机数C 为随机数避免 beta 路径被跳过 for (int i 0; i totalA; i) { hA[i] (float)rand() / RAND_MAX; }要注意随机数种子固定下来这样每次跑的结果可复现。有些人在每次运行时用time(NULL)做种子结果跑出来的性能波动还以为是自己优化出了问题其实只是数值变了导致里面的除法分支不同。固定种子是基准测试的规矩不是可有可无。4. 基准测试避坑指南为什么你的数字比别人低一半4.1 现象TFLOPS 只有理论峰值的 30%我第一次拿这套基准跑旧卡就是这样以为卡坏了后来发现是忘了开 Tensor Core。cuBLAS 默认的cublasSgemm在 FP32 下走的是普通 CUDA core 路径并不会自动用上 TF32 加速。要跑到更高性能得用cublasGemmEx配合CUBLAS_COMPUTE_32F_FAST_TF32。原因就是 cuBLAS 为了兼容性和精度把数学模式设成了默认的精确模式除非你显式放宽精度要求。解决方法是换用cublasGemmEx并把计算类型设为CUBLAS_COMPUTE_32F_FAST_TF32或者直接用 FP16 输入让内核自动走 Tensor Core。从那以后我跑基准之前都会先查一遍设备的计算能力确认哪种精度组合能吃满硬件。4.2 现象GEMV 带宽远低于显存带宽GEMV 的本质是访存密集型按理说带宽应该能打得很高。但很多人测出来连一半都不到。原因在于 A 矩阵是列主序存储的而 GEMV 的访问模式是逐行读取每个元素做乘加。如果trans参数设置不当相当于每次访存都跨越一整列局部性极差。解决方法是测两种转置模式CUBLAS_OP_N和CUBLAS_OP_T选带宽高的那个。另外当 M 和 K 都很小的时候GEMV 的启动开销占了大头合理做法是改用cublasSgemvBatched把多个向量一次性交给 GPU摊薄启动延迟。4.3 现象计时结果波动大同一配置两次跑差 20%这不是玄学是 GPU 的实时时钟频率在捣鬼。GPU 会根据负载动态升频和降频你跑基准时如果散热不好功耗墙就触发降频结果自然飘。解决方法是三管齐下一是开足--warmup次数让核心加热后进入稳定频率再计时二是增加迭代次数比如 200 次以上然后取最小值三是在运行前用nvidia-smi锁定时钟频率。锁时钟需要 root 权限如果没有至少要让测试程序反复运行同一个 kernel让频率稳定下来。4.4 现象换了一台 GPU 后相对性能排序完全颠倒同样的 cublas_benchmarks在 A 卡上 FP16 比 FP32 快 4 倍到 B 卡上只快 2 倍。原因是不同世代的硬件对 Tensor Core 的调度策略差异很大cuBLAS 内部维护着多个 kernel 变体根据尺寸、对齐和硬件代数自动选择。所以跨卡比较时不能只看 TFLOPS 高低还要把驱动的 CUDA 版本、cuBLAS 版本、卡的计算能力打出来否则你只能看到结果看不到背后的选路逻辑。我的习惯是把这些环境信息写进结果文件头部这样复现的时候不会一头雾水。4.5 现象第一次调用某个尺寸时特别慢后面就正常了这是 cuBLAS 的 lazy loading 机制在起作用。第一次调用某个 API 时它会加载对应 kernel 模块做各种 heuristics 搜索比如试探多个 tile 大小来选择最优配置这个开销可能高达几十毫秒。而正式的迭代循环只有几百微秒如果直接把第一次调用也计时最终平均时间会被严重拉高。解决方法是一定要做预热预热次数最好覆盖到每个不同的尺寸和转置组合。代码里可以设计成每个配置单独预热或者干脆跑一遍完整的空测试把所有 kernel 策略都触发一遍。5. 进阶让基准结果反映你的真实场景5.1 从 GEMM 扩展到 StridedBatch GEMM如果你的实际场景是小矩阵批量运算比如图神经网络里几百个邻接矩阵的乘法单测一个 GEMM 没有意义。cublas_benchmarks 这种框架很容易扩展出分块批量模式。常见做法是加一个--batch参数然后调用cublasStridedBatchedGemmEx。关键是要正确设置 stride 参数。对于连续存储的批量矩阵A 的 stride 是M*KB 的 stride 是K*NC 的 stride 是M*N。错一个 stride结果就是错的而且错得很有隐蔽性因为每个 batch 内部可能还算出了一个“合理”的数字。cublasStridedBatchedGemmEx(handle, transA, transB, m, n, k, alpha, dA, lda, strideA, dB, ldb, strideB, beta, dC, ldc, strideC, batchCount, computeType, CUBLAS_GEMM_DEFAULT_TENSOR_OP);参数里strideA、strideB、strideC都是元素个数不是字节数。这块最容易混。我吃过一次亏把字节数填进去结果前 256 个 batch 数据错了后面的因为内存越界直接报错。多看几遍文档或者用一个小 batch 尺寸验证数值能省下很多 debug 时间。批量模式下性能曲线和单发 GEMM 完全不同小尺寸时批量越大吞吐越高直到批量带来的调度开销超过并行收益。5.2 用 Roofline 模型解释基准数据拿到基准结果后我通常会先把每个点按计算强度分类画一条简单的手工 Roofline。不需要画图工具列表就行。对于 GEMM计算强度 2MNK / (MN NK MK)当 MNK4096 时强度大约为 1365 FLOPs/Byte远超机器平衡点属于计算密集。此时你做内存优化没用得靠提高计算密度比如换 Tensor Core。对于 GEMV计算强度只有 2K / (M*K ...) 约等于 2/M 当 M 很大时妥妥的访存密集。这时候锁频、调转置比什么都管用。基准代码里如果能输出计算强度整个性能分析会直观很多。5.3 自定义一个与你业务相同的测试配置标准基准测的是方阵和幂等尺寸但你的业务可能是 M1024、N128、K768 这种非对称尺寸。不用改代码只要在调用入口传入这些值即可。我见过有人为了测某个深度优化的卷积等效 GEMM直接把m/n/k设成卷积输出的形状参数这样能提前判断 cuBLAS 的通用实现是否够用还是必须自己写 kernel。如果你想测得更细可以在BenchConfig里加一个alpha和beta的字段。很多人默认 alpha1、beta0但真实场景里 beta 往往是 1也就是累加操作。cuBLAS 对 beta0 和 beta1 可能会走不同的 kernel 路径beta1 时可能要多读一次 C 再写回性能差几个百分点。5.4 验证基准本身没有测错最后一个要养成的习惯是每次新增测试场景时先跑一遍带--check的验证模式确认结果数值正确再开始谈性能。我过去有一次为了追求速度把迭代次数加到 2000结果矩阵尺寸太大导致数据溢出TFLOPS 看起来很高实际算出的 C 已经全是无穷大。你要是在这种错误数据上做优化后面所有优化决策都会跟着错。验证模式的做法很简单在 CPU 上用双精度跑一遍同样的运算和 GPU 结果比较相对误差小于 1e-3 就通过。从那以后我每次做性能优化都会强制先跑一遍基准并验证数值确认环境没问题再看性能曲线。这套流程虽然多花两分钟但能省下后面一小时的 debug 时间。希望帮到你。本文还有配套的精品资源点击获取
