cuBLAS 与 cuDNN：为什么生产环境的 GPU 代码都调库
更新时间：2026-09-25。矩阵乘法 tiling 一篇里，我们亲手把一个 GEMM kernel 从几个 GFLOPS 优化到几百个 GFLOPS——但那只是入门级的分块核，距离峰值算力仍有数量级差距。本篇回答一个自然的追问：既然道理都懂了，为什么生产代码里几乎看不到人手写 GEMM 和卷积，全部转而调用 cuBLAS、cuDNN 这样的闭源库？底层硬件为什么值得这样封装，见 Tensor Core 张量核心；库之上框架如何再做算子融合，见 算子融合与编译。
本文要回答的问题
- 一个 GEMM 或卷积 kernel「写对」和「写到峰值」之间差多少工作量？
- 库里的 autotuning 到底在做什么？它凭什么按架构、形状、精度自动选 kernel？
- cuBLAS 的核心对象是什么（handle、列主序、stream），调用一次 GEMM 要注意什么？
- cuDNN 的 descriptor 模型怎么描述卷积？workspace 是什么、为什么要自己分配？
- cuSPARSE、cuFFT 各自解决什么问题，什么场景下选哪个库？
- 实战中有哪些高频坑：布局转置、首次 autotune 延迟、版本与 GPU 架构匹配？
一、写对 kernel 与写到峰值之间，差几百人月
先给「自己写」算一笔成本账。教学版 tiled GEMM 大约一百行，能跑到峰值的 5%~10%；要继续逼近峰值，需要依次解决：shared memory 的 swizzle 布局消除 bank conflict、cp.async 双缓冲把搬运藏进计算、寄存器分片与累加器驻留、warp 间的 L 形分块（split-K）、Tensor Core 的 fragment 装载与 K 维对齐、尾部非整除维度的 mask 处理。每一项都不是写出来就行，而要在具体卡上反复 profile、反复折中：tile 调大了 occupancy 掉，调小了阵列喂不饱。
卷积的复杂度更高。同一个卷积，在不同形状下最快的算法完全不同：小通道用隐 GEMM（implicit GEMM），大特征图用 Winograd 把乘法数量压低，1×1 卷积直接退化成 GEMM，深度可分离卷积又是另一套访存模式。每种算法还要再乘上数据类型（FP32/TF32/FP16/BF16/FP8）、布局（NCHW/NHWC）、微架构（Volta/Ampere/Hopper/Blackwell 各代 Tensor Core 指令不同）的组合。组合爆炸的结果是：一个覆盖主流形状、各档精度、多代架构的高性能 GEMM/卷积库，背后是厂商团队数百人年的持续投入，而且每出新架构就要再来一轮。
这笔账决定了行业分工：NVIDIA 把硬件细节和调优知识封进 cuBLAS、cuDNN 等库，随 CUDA 版本免费发布；应用团队只描述「要算什么」，不描述「怎么算」。一个直接可验证的收益是：在形状足够大（如 M=N=K=4096 以上）、维度对齐、精度匹配的 GEMM 上，cuBLAS 稳定达到对应硬件峰值算力的 90% 以上，大矩阵常落在 90%~95% 区间；而同等形状下即便是熟练工程师手写的 kernel，通常只在峰值的 30%~70%。形状越小、精度越新，这个差距越大。
二、库分层：你的代码和硬件之间隔着什么
下图是一次典型深度学习计算的调用栈。框架和业务代码不直接触碰 kernel，全部计算下沉到几家「专算子库」；库里装着为每代架构预编译的几百个 kernel 变体，运行时才决定执行哪一个。
autotuning 是库价值的核心机制，它分两个层次。第一层是静态启发式选择：调用 GEMM 时，库拿到矩阵维度 M、N、K、数据类型、转置标志和当前 GPU 的计算能力编号（compute capability），查表式地从候选 kernel 中挑一个——这个过程很快，但不实测。第二层是实测式 autotuning：cuDNN 提供 cudnnFindConvolutionForwardAlgorithm 等接口，让框架把若干候选算法真跑一遍（通常是限制迭代次数的小样本计时），按实测耗时排序，框架随后缓存「形状+精度→最优算法」的映射。下图给出选择依据。
理解这张图，就理解了库的本质：它卖的不是算子实现，而是「把几百人年的调优知识压缩成一次查表」的能力。代价则是抽象层必然带来的约束——你得按它的方式描述问题，这正是后面几节的内容。
三、cuBLAS：BLAS 三级、handle、列主序与 stream
cuBLAS 是标准 BLAS（Basic Linear Algebra Subprograms）的 GPU 实现。BLAS 按计算层次分三级：Level 1 是向量-向量运算（如 axpy，O(n) 计算 O(n) 数据），Level 2 是矩阵-向量运算（如 gemv，O(n²) 计算 O(n²) 数据），Level 3 是矩阵-矩阵运算（如 gemm，O(n³) 计算 O(n²) 数据）。性能故事几乎全在 Level 3：算术强度随矩阵维度线性增长，只有它能把 Tensor Core 喂饱。cuBLAS 的 cublasGemmEx（混合精度扩展版）就是深度学习里全连接层、注意力投影、大矩阵乘的最终落点。
理解 GEMM 今天的地位，要看 Transformer 的计算构成。一个大语言模型的训练 FLOPs 里，GEMM 类运算占 90% 以上：自注意力中的 Q/K/V 三个输入投影、输出投影，以及前馈网络（FFN）的两个大线性层，全部是 GEMM；连注意力分数本身也是两次批 GEMM——Q·K^T 和注意力权重乘 V，只有中间的 softmax 不是。batch 和序列长度一大，这些 GEMM 的 M、N、K 动辄数千，正是库最能发挥的形状。也正因如此，业界对 GPU 的性能讨论几乎等价于对 GEMM 性能的讨论，cuBLAS 之上还衍生出 cuBLASLt，专门增加灵活的数据布局（如行列交错的矩阵排布）、批量广播和 epilogue 融合（乘加后直接接 bias、ReLU、GELU），PyTorch 的 addmm、linear 在现代版本上大多走的是 cuBLASLt 而非传统 cuBLAS 入口。
使用模型有四个必须建立的概念：
handle（句柄）：cublasCreate(&handle) 创建，它承载库的上下文并指向当前设备。handle 应在初始化时创建一次、全程复用，而不是每次调用都建——句柄内部关联了 stream 与内部状态，反复创建是实测可见的性能浪费。线程安全方面要记住一条规则：同一个 handle 被多个线程并发使用是不安全的，库内部的 stream 指针和工作区状态会互相踩踏；而多个线程各自持有独立 handle 则完全没问题。因此多线程推理服务的标准模式是每个线程（或每条 stream）一个 handle，初始化时建好、线程内复用。还要注意 cublasSetStream 改的是这个 handle 内部绑定的 stream，handle 本身不携带设备所有权，但创建时所在设备会被记住——跨设备使用前要先 cudaSetDevice，否则调用会落到错误的 GPU 上。
列主序（column-major）：BLAS 沿袭 Fortran 传统，矩阵按列连续存储；而 C/C++ 与多数框架的张量是行主序。直接对接有两种标准做法：一是利用「行主序矩阵的 C = A·B，等价于列主序下 C^T = B^T·A^T」，交换 A、B 的顺序并互换转置参数；二是把数据按库期望的布局摆放。这是 cuBLAS 使用中出错率最高的点，没有之一。
转置与 leading dimension：cublasGemmEx 的签名是 C = alpha·op(A)·op(B) + beta·C，其中 op 由 CUBLAS_OP_T/N 决定是否转置，lda/ldb/ldc 是主维跨度。行列主序换算时 lda 该填什么必须按列主序语义重新推导，照抄数学维度是典型错误。
stream 接入：cublasSetStream(handle, stream) 让后续调用提交到指定 CUDA stream，从而与 memcpy、其他 kernel 异步并发。多线程共享 GPU 时，通常每线程一个 handle、各自绑定自己的 stream，避免调用互相串行化。
一段最小骨架大致是：创建 handle → cublasSetStream → 准备设备指针 → 调 cublasGemmEx（指定 CUDA_R_16F 计算类型与 CUBLAS_COMPUTE_32F，即 FP16 输入、FP32 累加）→ 同步或交给事件管理。整个过程中没有一行 kernel 代码，但 Tensor Core、swizzle、双缓冲全部在库内部生效。
四、cuDNN：descriptor 模型与 workspace
cuDNN 面向神经网络算子：卷积、池化（pooling）、激活（activation，如 ReLU）、归一化、RNN/LSTM 以及它们的前后向。它与 cuBLAS 的抽象风格不同——cuBLAS 用参数列表描述矩阵，cuDNN 用 descriptor（描述符） 描述一切张量和算子，典型调用流程是：
cudnnCreate(&cudnn)创建库上下文；- 用
cudnnSetTensor4dDescriptor（或通用的nd版本）分别为输入、输出、权重建张量描述符，写清数据类型、布局（NCHW/NHWC）和各维尺寸； - 用
cudnnSetConvolution2dDescriptor建卷积描述符，写清 padding、stride、dilation、卷积模式与累加精度； - 调
cudnnGetConvolutionForwardAlgorithm_v7或 find 系列拿到算法； - 查询所需 workspace 大小，自行
cudaMalloc一块显存传给cudnnConvolutionForward。
workspace（工作区） 是 cuDNN 特有的概念。高性能卷积算法（implicit GEMM、Winograd、FFT 类）需要临时显存存放变换中间结果；库不替你管理这块内存，而是「查询大小、调用方分配、传入指针」。这样设计的好处是应用可以统一规划显存池、在多个算子间复用同一块 scratch，坏处是新手常忘记 workspace 也计入峰值显存——OOM 往往不是张量太大，而是给几个并发卷积各留了一份 workspace。
第一节提到卷积有多种算法，这里把三条主流路线和它们的 workspace 特征讲具体。implicit GEMM（隐式 GEMM） 把卷积等价为「im2col 展开后的大矩阵乘权重矩阵」，但不真的把展开矩阵写进显存，而是在 GEMM 读数据时现场算出每个元素对应输入的哪个位置——通用性最好，workspace 通常为零或只有很小一块（几 MB 量级），现代 cuDNN 在大多数形状上的默认选择就是它。Winograd 对 3×3 卷积特别有效：以 F(2×2, 3×3) 为例，算一个 2×2 输出块只需 4 次乘法而不是朴素的 36 次，F(4×4, 3×3) 还能进一步压低；代价是输入、输出都要做域变换，变换结果要落在 workspace 里，大 batch、大通道下 workspace 可达几十到上百 MB，且变换引入的数值误差在小核卷积之外可能变得明显，FP16 下需留意精度。FFT 把卷积变到频域做点乘，对大核（5×5、7×7）收益可观，但变换和复数中间结果的 workspace 常以数百 MB 计，小核上反而吃亏，现代 GPU 上已很少成为默认。每种算法在具体形状下的 workspace 大小都要用 cudnnGetConvolutionForwardWorkspaceSize 实测查询，不能凭经验估。
API 层面还要知道 cuDNN 8 引入的 frontend（graph）API。前面描述的 Set...Descriptor 加逐算子调用是传统用法；cuDNN 8 起官方主推图式接口：把卷积、偏置、加法、激活等表示成一张 operation graph，交给库的引擎启发式选择 execution plan，原生支持 conv+bias+ReLU 这类融合，并能把 plan 序列化落盘、跨进程复用（即前面预热一节说的 engine caching）。新代码建议直接用 frontend 头文件库，旧的逐算子 API 仍可用但不再获得融合能力上的增强。
与 cuBLAS 一样，cuDNN 在满足条件时自动走 Tensor Core：数据类型选 FP16/BF16/FP8、配置匹配（CUDNN_TENSOR_OP_MATH 相关开关、NHWC 布局在多数代际上对 Tensor Core 更友好）后，卷积也能逼近峰值。反过来，NCHW 布局配 FP16 在部分旧架构上会悄悄退回非 Tensor Core 路径，结果完全正确、速度差数倍——这种「静默不加速」只能靠 profile 发现。
给一组可建立量级感的数字。A100 的稠密 FP16/BF16 Tensor Core 峰值是 312 TFLOPS（2:4 结构化稀疏下标称 624，但稀疏需要模型权重满足模式，并非白拿），大形状 FP16 GEMM 在 cuBLAS 里实测常落在 280 TFLOPS 以上，即峰值的九成；V100 的稠密 FP16 峰值为 125 TFLOPS，同口径大 GEMM 约 110 TFLOPS。卷积的利用率略低，形状规整的主流 CNN 前向通常在峰值的 50%~80%，Winograd/implicit GEMM 命中良好时接近 GEMM 水平。而端到端训练因为还包含 softmax、归一化、optimizer、通信与大量小算子，业界报告的 MFU（Model FLOPs Utilization）普遍只有峰值的 30%~50%——单算子 90% 与整模型四成之间的差距，正是融合、编译和并行策略要追回的部分。
五、cuSPARSE、cuFFT 与选型对照表
GPU 的算子库不止两家。其余常见成员还包括：cuSPARSE 处理稀疏矩阵（CSR/COO/CSC 格式的 SpMV、SpMM、稀疏三角求解），用于图计算、稀疏线性求解和结构化剪枝网络；cuFFT 提供批量 FFT，用于信号处理、频域卷积和部分科学计算；cuRAND 做随机数生成；cuSOLVER 做稠密/稀疏分解与特征值；TensorRT 则进一步把推理图整体编译优化（融合、量化、kernel 选择），属于更高一层的部署库。
| 库 | 核心对象/模型 | 主要算子 | 何时用 |
|---|---|---|---|
| cuBLAS | handle + 列主序矩阵参数 | GEMM、GEMV、axpy 等三级 BLAS | 全连接、注意力投影、任何稠密矩阵乘 |
| cuDNN | 库句柄 + tensor/conv descriptor + workspace | 卷积、池化、激活、归一化、RNN 前后向 | CNN 与神经网络专用算子 |
| cuSPARSE | sparse descriptor + CSR/COO 格式矩阵 | SpMV、SpMM、稀疏求解 | 稀疏度高的矩阵、图计算、剪枝推理 |
| cuFFT | plan（含维数与批量） | 正/逆 FFT、批量 FFT | 频谱分析、频域滤波、FFT 卷积 |
| cuSOLVER | handle + 矩阵描述符 | LU/QR/Cholesky 分解、特征值 | 需要矩阵分解的科学/优化计算 |
| TensorRT | 网络定义/引擎构建 | 整图融合、精度校准、kernel autotune | 训练后的推理部署，追求端到端延迟 |
选型的基本判据是算子的数据形态：稠密规整矩阵乘找 cuBLAS，神经网络模式化算子找 cuDNN，稀疏结构找 cuSPARSE，频域找 cuFFT。而当多个小算子串在一起、库调用之间的 launch 开销和显存往返成为瓶颈时，下一步不是再换库，而是上融合与编译，参见 算子融合与编译。
六、实战中的四个高频坑
布局与转置错误：症状是结果矩阵数值全错或形状错位。根因永远是行/列主序与 OP_T、leading dimension 的组合没对上。排查方法固定：先用小矩阵（如 2×3、3×4）单测，与 CPU 结果逐元素对拍，再上大形状。
workspace 与显存规划：除了上文说的显存占用，还要注意 find 阶段可能为候选算法临时分配多份 workspace，显存紧张时 autotune 本身可能 OOM。生产做法是限定候选算法数量、复用显存池，并把「选定算法 + workspace 大小」固化下来。
首次 autotune 延迟：服务上线第一批请求特别慢，原因是 find/启发式初始化和 kernel JIT（新库版本对特定形状可能现场编译）。推理服务必须预热：上线前用真实形状跑 dummy 推理填充缓存，或用 cuDNN 的持久化计划（cudnn frontend + engine caching）把选择结果落盘，跨进程复用。
版本与架构匹配：库必须「认识」你的 GPU 才能调对 kernel——在 Hopper 上用 CUDA 11 时代的旧 cuBLAS，要么用不上新 Tensor Core 指令，要么直接不支持；而容器里框架自带的 cuBLAS 与系统驱动支持的最高 CUDA 版本不一致，是环境类故障的第一来源。部署前核对三件套：GPU 架构 compute capability、驱动支持的 CUDA 版本、框架链接的 cuBLAS/cuDNN 版本。额外提醒，cuDNN 主版本升级偶尔伴随算法选择和默认精度行为变化，升级后要重跑性能基线，不能只验证正确性。
最后建立一个合理预期：库保证的是「大形状、规整、命中支持路径」时的高水平下限，不是所有形状都快。M、N、K 很小或某一维极端窄小时（典型如自回归解码的 batch=1 GEMM），kernel launch 开销和边角浪费占主导，库也只能做到峰值的几个百分点——这不是库的问题，而是算子形状决定的，此时应从 batch 聚合、矩阵合并等算法层面求解。
这一篇在知识体系里的位置
- 软件手法上游：矩阵乘法 tiling 讲库内部 kernel 赖以加速的分块原理；
- 硬件落点：Tensor Core 张量核心 讲库自动调用的专用矩阵单元与各精度算力；
- 继续向上：算子融合与编译 讲库调用之上，框架如何消除 launch 与访存开销。
总结
生产代码调库而不自己写 kernel，是因为一个覆盖多代架构、各档精度与形状组合的 GEMM/卷积实现要花厂商数百人年调优，而 cuBLAS、cuDNN 把这些知识压缩成运行时一次 autotuning：cuBLAS 以 handle、列主序、转置参数和 stream 承接三级 BLAS 的 GEMM，cuDNN 以 tensor/conv descriptor 加调用方自管的 workspace 承接卷积、池化、激活与 RNN，cuSPARSE、cuFFT、cuSOLVER、TensorRT 则按稀疏、频域、分解、整图部署各自分工；命中大形状与对齐路径时 cuBLAS GEMM 稳定达到峰值 90% 以上，实战中要防的是行列主序转置错误、workspace 显存漏算、首次 autotune 延迟需预热，以及库版本与 GPU 架构必须匹配——形状窄小的低利用率则是算法形状问题，调库解决不了，得从 batch 聚合层面求解。