简介：cublas_benchmarks 是一份面向 CUDA 高性能计算开发者与深度学习底层优化人员的 C++ 基准测试资源，聚焦 NVIDIA CuBLAS 库在 GEMM（通用矩阵乘法）与 GEMV（矩阵向量乘法）两类核心运算上的性能表现。资源包共 10 个文件，以 4 个 .cu 与 4 个 .h 头文件为主体，分别承载矩阵/向量数据结构、CuBLAS 函数封装与性能计时逻辑，另附 1 个 Makefile 与 1 个 README，压缩包仅 13KB，结构紧凑、便于快速编译运行。已有 1095 人学习下载。读者可借助其中的 xgemm、xgemv 测试代码，调整 M、N、K 等矩阵规模参数，测量不同硬件与 CUDA 版本下的执行耗时，进而理解 GPU 加速线性代数的性能瓶颈，为代码优化与计算策略选型提供实测参考。

1. 从一次 kernel 调优翻车说起：这个基准项目到底测什么

有个做推理加速的朋友，接手一个矩阵乘优化任务，拍脑袋把 tile 从 128 调到 256，本地跑出来快了 8%，结果上线后整体吞吐反而掉了。问题出在他只测了单个 shape，没覆盖真实业务里那些“瘦长”矩阵。后来他翻出一个专门测 GEMM 和 GEMV 的基准项目，把常见 shape 全跑了一遍，才发现自己那个调参只对特定尺寸有效。这类基准项目就是干这个的——它不训练模型、不做推理，只专注一件事：把 cuBLAS 里 GEMM（通用矩阵乘）和 GEMV（矩阵向量乘）在不同规模、不同数据类型下的性能摸清楚。适合谁？写 CUDA kernel 的、做算子优化的、调推理框架底层的人。你不需要它来跑业务，但你需要它来回答“这个 shape 到底该用哪个 API、哪个 tile、哪个精度”。C++ 在这里是主力语言，因为要直接调 CUDA Runtime 和 cuBLAS，还要做高精度计时。

2. 项目结构拆解：从 CMake 到 benchmark 入口的完整链路

2.1 目录布局与构建系统

这类基准项目通常不会太复杂，但结构必须清晰，否则你加一个测试用例就要改三处。我拿到源码后第一件事是看目录，典型布局长这样：

cublas_benchmarks/ ├── CMakeLists.txt ├── src/ │ ├── main.cpp │ ├── benchmark_runner.cpp │ ├── benchmark_runner.h │ ├── gemm_bench.cpp │ ├── gemv_bench.cpp │ └── utils/ │ ├── timer.h │ ├── timer.cpp │ └── device_info.cpp ├── include/ │ └── cublas_bench/ │ └── types.h └── scripts/ └── run_all.sh

CMakeLists.txt是入口，核心就三件事：找 CUDA、找 cuBLAS、把源文件编成可执行文件。常见写法是：

cmake_minimum_required(VERSION 3.18) project(cublas_benchmarks LANGUAGES CXX CUDA) set(CMAKE_CXX_STANDARD 17) set(CMAKE_CXX_STANDARD_REQUIRED ON) find_package(CUDAToolkit REQUIRED) add_executable(cublas_bench src/main.cpp src/benchmark_runner.cpp src/gemm_bench.cpp src/gemv_bench.cpp src/utils/timer.cpp src/utils/device_info.cpp ) target_include_directories(cublas_bench PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/include ) target_link_libraries(cublas_bench PRIVATE CUDA::cudart CUDA::cublas )

这里find_package(CUDAToolkit REQUIRED)是 CMake 3.17 之后的标准做法，比手动find_library干净得多。CUDA::cublas会自动带上头文件路径和链接库，不用你写-lcublas。注意LANGUAGES CXX CUDA里带了 CUDA，因为后面可能混编.cu文件；如果全是.cpp调 cuBLAS API，其实只写 CXX 也行，但加上 CUDA 更保险。

2.2 基准入口与参数解析

main.cpp通常很薄，只做参数解析和分发。我一般会设计成子命令模式，比如./cublas_bench gemm --m 1024 --n 1024 --k 1024这样。核心逻辑在benchmark_runner里：

// benchmark_runner.h struct BenchmarkConfig { std::string op; // "gemm" or "gemv" int m, n, k; int warmup_iters; int bench_iters; cublasOperation_t transa; cublasOperation_t transb; cudaDataType_t data_type; }; class BenchmarkRunner { public: explicit BenchmarkRunner(const BenchmarkConfig& cfg); void run(); void print_result() const; private: BenchmarkConfig cfg_; float elapsed_ms_; double gflops_; };

warmup_iters和bench_iters是两个关键参数。warmup 是为了让 GPU 频率稳定、cache 预热，通常设 10 到 50 次；bench 是实际计时次数，一般 100 次起步。少了这两个，测出来的数就是玄学。

2.3 计时器实现：为什么不能用 CPU 时钟

utils/timer.h是整个基准的“尺子”，尺子不准后面全白搭。常见做法是用 CUDA Event：

// timer.h class CudaTimer { public: CudaTimer() { cudaEventCreate(&start_); cudaEventCreate(&stop_); } ~CudaTimer() { cudaEventDestroy(start_); cudaEventDestroy(stop_); } void start() { cudaEventRecord(start_, 0); } void stop() { cudaEventRecord(stop_, 0); } float elapsed_ms() { float ms = 0.0f; cudaEventSynchronize(stop_); cudaEventElapsedTime(&ms, start_, stop_); return ms; } private: cudaEvent_t start_, stop_; };

cudaEventRecord把事件插到流里，cudaEventElapsedTime返回两个事件之间的毫秒数。精度通常在 0.5 微秒左右，比std::chrono测 CPU 时间靠谱得多，因为 CPU 时间会把 kernel launch 开销和异步执行混在一起。注意cudaEventSynchronize(stop_)不能省，否则你读到的是还没执行完的时间。

3. 核心模块拆解：GEMM 与 GEMV 的基准代码逐段讲

3.1 GEMM 基准：从内存分配到 kernel 调用

GEMM 测的是C = alpha * A * B + beta * C，cuBLAS 里对应cublasSgemm（单精度）或cublasGemmEx（混合精度）。一个完整的基准函数长这样：

// gemm_bench.cpp void run_gemm_bench(const BenchmarkConfig& cfg) { // 1. 分配 host 和 device 内存 size_t size_a = (size_t)cfg.m * cfg.k * sizeof(float); size_t size_b = (size_t)cfg.k * cfg.n * sizeof(float); size_t size_c = (size_t)cfg.m * cfg.n * sizeof(float); float *h_a = (float*)malloc(size_a); float *h_b = (float*)malloc(size_b); float *h_c = (float*)malloc(size_c); float *d_a, *d_b, *d_c; cudaMalloc(&d_a, size_a); cudaMalloc(&d_b, size_b); cudaMalloc(&d_c, size_c); // 2. 初始化数据（用随机数避免全零导致优化） srand(42); for (size_t i = 0; i < size_a / sizeof(float); ++i) h_a[i] = (float)rand() / RAND_MAX; for (size_t i = 0; i < size_b / sizeof(float); ++i) h_b[i] = (float)rand() / RAND_MAX; // 3. 拷贝到 device cudaMemcpy(d_a, h_a, size_a, cudaMemcpyHostToDevice); cudaMemcpy(d_b, h_b, size_b, cudaMemcpyHostToDevice); // 4. 创建 cuBLAS handle cublasHandle_t handle; cublasCreate(&handle); // 5. warmup float alpha = 1.0f, beta = 0.0f; for (int i = 0; i < cfg.warmup_iters; ++i) { cublasSgemm(handle, cfg.transa, cfg.transb, cfg.m, cfg.n, cfg.k, &alpha, d_a, cfg.m, d_b, cfg.k, &beta, d_c, cfg.m); } cudaDeviceSynchronize(); // 6. 计时 CudaTimer timer; timer.start(); for (int i = 0; i < cfg.bench_iters; ++i) { cublasSgemm(handle, cfg.transa, cfg.transb, cfg.m, cfg.n, cfg.k, &alpha, d_a, cfg.m, d_b, cfg.k, &beta, d_c, cfg.m); } timer.stop(); float elapsed = timer.elapsed_ms(); // 7. 计算 GFLOPS double flops = 2.0 * cfg.m * cfg.n * cfg.k * cfg.bench_iters; double gflops = flops / (elapsed * 1e6); // 8. 清理 cublasDestroy(handle); cudaFree(d_a); cudaFree(d_b); cudaFree(d_c); free(h_a); free(h_b); free(h_c); printf("GEMM m=%d n=%d k=%d: %.3f ms, %.2f GFLOPS\n", cfg.m, cfg.n, cfg.k, elapsed / cfg.bench_iters, gflops); }

逐段说。第 1 步分配内存时，size_a用(size_t)强转是为了防止m * k在 int 下溢出——大矩阵很容易超 2^31。第 2 步用随机数而不是全零，因为 cuBLAS 内部可能对零矩阵做短路优化，测出来偏快。第 3 步的cudaMemcpy是同步的，确保数据到位再开跑。第 4 步cublasCreate会初始化 cuBLAS 上下文，第一次调用有额外开销，所以必须放在 warmup 之前。第 5 步 warmup 循环里不计时，但每次调用后要cudaDeviceSynchronize()确保队列排空。第 6 步计时循环里连续发 kernel，最后统一同步，这样测的是纯 kernel 执行时间。第 7 步 GFLOPS 公式里2.0 * m * n * k是因为一次乘加算两个浮点操作。第 8 步清理顺序无所谓，但别漏。

3.2 GEMV 基准：内存带宽才是瓶颈

GEMV 是y = alpha * A * x + beta * y，对应cublasSgemv。它和 GEMM 最大的区别是：GEMV 是内存带宽受限，不是计算受限。所以基准代码里要特别关注数据搬运量：

// gemv_bench.cpp void run_gemv_bench(const BenchmarkConfig& cfg) { // A: m x n, x: n x 1, y: m x 1 size_t size_a = (size_t)cfg.m * cfg.n * sizeof(float); size_t size_x = (size_t)cfg.n * sizeof(float); size_t size_y = (size_t)cfg.m * sizeof(float); float *d_a, *d_x, *d_y; cudaMalloc(&d_a, size_a); cudaMalloc(&d_x, size_x); cudaMalloc(&d_y, size_y); // 用 cudaMemset 快速初始化，GEMV 对数值不敏感 cudaMemset(d_a, 0, size_a); cudaMemset(d_x, 0, size_x); cudaMemset(d_y, 0, size_y); cublasHandle_t handle; cublasCreate(&handle); float alpha = 1.0f, beta = 0.0f; // warmup for (int i = 0; i < cfg.warmup_iters; ++i) { cublasSgemv(handle, CUBLAS_OP_N, cfg.m, cfg.n, &alpha, d_a, cfg.m, d_x, 1, &beta, d_y, 1); } cudaDeviceSynchronize(); CudaTimer timer; timer.start(); for (int i = 0; i < cfg.bench_iters; ++i) { cublasSgemv(handle, CUBLAS_OP_N, cfg.m, cfg.n, &alpha, d_a, cfg.m, d_x, 1, &beta, d_y, 1); } timer.stop(); float elapsed = timer.elapsed_ms(); // GEMV 看有效带宽：读 A + 读 x + 写 y double bytes = (double)size_a + size_x + size_y; double bandwidth = bytes * cfg.bench_iters / (elapsed * 1e-3) / 1e9; printf("GEMV m=%d n=%d: %.3f ms, %.2f GB/s\n", cfg.m, cfg.n, elapsed / cfg.bench_iters, bandwidth); cublasDestroy(handle); cudaFree(d_a); cudaFree(d_x); cudaFree(d_y); }

注意cublasSgemv的lda参数传的是cfg.m，因为 cuBLAS 是列主序，A 是 m x n 但按列存，leading dimension 是 m。incx和incy都传 1，表示连续存储。带宽公式里bytes只算了一次读 A、读 x、写 y，实际还有 y 的读（beta 不为零时），但 beta=0 时可以忽略。这个数直接告诉你 kernel 有没有打满显存带宽——如果测出来只有理论带宽的 30%，那说明访问模式有问题。

3.3 批量测试与结果输出

单点测试不够，真实场景要扫一批 shape。常见做法是在main.cpp里读一个 shape 列表，循环调用：

// main.cpp 片段 std::vector<BenchmarkConfig> configs = { {"gemm", 256, 256, 256, 10, 100, CUBLAS_OP_N, CUBLAS_OP_N, CUDA_R_32F}, {"gemm", 512, 512, 512, 10, 100, CUBLAS_OP_N, CUBLAS_OP_N, CUDA_R_32F}, {"gemm", 1024, 1024, 1024, 10, 100, CUBLAS_OP_N, CUBLAS_OP_N, CUDA_R_32F}, {"gemm", 2048, 2048, 2048, 10, 50, CUBLAS_OP_N, CUBLAS_OP_N, CUDA_R_32F}, {"gemv", 4096, 4096, 0, 10, 100, CUBLAS_OP_N, CUBLAS_OP_N, CUDA_R_32F}, }; for (const auto& cfg : configs) { BenchmarkRunner runner(cfg); runner.run(); runner.print_result(); }

输出建议同时写 CSV，方便后面画图：

// print_result 里追加 FILE* fp = fopen("results.csv", "a"); fprintf(fp, "%s,%d,%d,%d,%.3f,%.2f\n", cfg_.op.c_str(), cfg_.m, cfg_.n, cfg_.k, elapsed_ms_, gflops_); fclose(fp);

这样跑完一轮，results.csv直接丢进 Python 画曲线，比在终端里瞪眼看数字高效得多。

4. 避坑与排查：基准测试里最容易翻车的五个点

4.1 现象：测出来比理论峰值还高

原因通常是没做 warmup，第一次调用 cuBLAS 时内部有初始化开销，但如果你把第一次的耗时也算进平均，反而可能因为 GPU 还没降频而偏快；更常见的是计时循环里没同步，kernel 还在队列里就停了表。解决：warmup 至少 10 次，计时循环结束后统一cudaDeviceSynchronize()再读时间。

4.2 现象：同一 shape 跑三次结果差 20%

原因一般是 GPU 频率波动。数据中心卡有固定频率还好，消费级卡会动态调频。解决：跑之前用nvidia-smi -lgc锁频（需要权限），或者至少跑 5 轮取中位数。另一个可能是显存 ECC 开了，纠错会拖慢带宽，查nvidia-smi -q | grep ECC。

4.3 现象：GEMV 带宽只有理论值的三分之一

原因多半是lda传错了。cuBLAS 列主序下，A 是 m x n，leading dimension 是 m 不是 n。如果你按行主序的习惯传了 n，cuBLAS 会按错误的步长读内存，cache 命中率暴跌。解决：确认cublasSgemv的lda参数等于矩阵行数（列主序下）。

4.4 现象：大矩阵报错cudaErrorMemoryAllocation

原因：m * n * sizeof(float)在 int 下溢出成负数，cudaMalloc收到一个巨大的 size_t。解决：所有尺寸计算强制用size_t或long long，并且检查cudaMalloc返回值。另外，如果显存不够，可以分块测或者用cudaMallocManaged走统一内存，但后者会引入页迁移开销，测出来的数不能直接当 kernel 性能。

4.5 现象：换一台机器结果完全对不上

原因：不同 GPU 架构的 SM 数量、显存带宽、L2 cache 大小都不一样。解决：基准结果必须带设备信息。在device_info.cpp里用cudaGetDeviceProperties打印设备名、SM 数、显存带宽、L2 大小，输出到 CSV 第一行。没有设备信息的 benchmark 结果就是黑匣子，换台机器就没法比。

5. 二次开发：加自定义 shape 扫描与精度对比的实操

5.1 扩展 shape 生成器

原项目如果只支持手动列 shape，你二次开发第一件事就是加自动扫描。我一般会写一个生成器，按对数间隔采样：

std::vector<int> generate_sizes(int min_size, int max_size, int num_points) { std::vector<int> sizes; double log_min = std::log(min_size); double log_max = std::log(max_size); for (int i = 0; i < num_points; ++i) { double t = (double)i / (num_points - 1); int s = (int)std::exp(log_min + t * (log_max - log_min)); // 对齐到 32 的倍数，避免非对齐访问干扰 s = (s + 31) / 32 * 32; sizes.push_back(s); } return sizes; }

这样从 128 到 8192 采 20 个点，覆盖小、中、大矩阵。对齐到 32 是因为 cuBLAS 内部 tile 通常按 32 对齐，非对齐 shape 会触发边界处理，测出来偏慢但不代表真实场景。

5.2 加混合精度对比

现在推理场景 FP16 和 BF16 很常见，你可以在BenchmarkConfig里加data_type字段，然后分派到cublasGemmEx：

if (cfg.data_type == CUDA_R_16F) { cublasGemmEx(handle, cfg.transa, cfg.transb, cfg.m, cfg.n, cfg.k, &alpha, d_a, CUDA_R_16F, cfg.m, d_b, CUDA_R_16F, cfg.k, &beta, d_c, CUDA_R_16F, cfg.m, CUBLAS_COMPUTE_32F, // 累加用 FP32 CUBLAS_GEMM_DEFAULT_TENSOR_OP); }

关键参数是CUBLAS_COMPUTE_32F，表示乘加在 FP32 下累加，输入输出是 FP16。这样测出来的精度和速度都接近真实推理。注意CUBLAS_GEMM_DEFAULT_TENSOR_OP会启用 Tensor Core，如果 shape 不满足对齐要求可能回退到普通 kernel，测出来会跳变。

5.3 验证方法：用 cuBLAS 官方文档对表

改完之后怎么确认没测错？我的习惯是拿一个已知 shape 去对 cuBLAS 官方文档里的性能表。比如 4096x4096x4096 FP32，在某个架构上官方标称多少 TFLOPS，你测出来应该在 80% 到 110% 之间。低于 80% 说明你的计时或 warmup 有问题，高于 110% 说明你可能漏算了某些开销。另一个办法是跟cublasLt的 heuristic 结果对比，如果差太多，检查是不是没开 Tensor Core。

5.4 一个我踩过的坑

有次我加了一个--transa T的选项，想测转置矩阵的性能，结果忘了在cublasSgemm里同步改lda。转置后 A 的 leading dimension 应该变成 k 而不是 m，我没改，测出来的数比不转置还快——因为 cuBLAS 读错了内存，实际算的是另一个矩阵。从那以后我每次加新参数，都强制走一遍“小矩阵手算验证”：用 4x4 的矩阵，CPU 算一遍结果，GPU 算一遍，对上了再跑大矩阵。这个习惯帮我省了至少三次返工。希望帮到你。

本文还有配套的精品资源，点击获取