搞模型部署的迟早要跟 CUDA 打交道。不管你用的是 TensorRT、vLLM、ONNX Runtime 还是自己写推理脚本跑一段时间一定会撞上这些问题怎么这么慢、显存怎么又不够了、这个算子为什么死活打不满 GPU、报错信息里那一串看不懂的 “sm_120” 到底是什么意思。这时候回头去补 CUDA 编程和 GPU 架构的课不是学术需求是实打实的生产需求。这个系列前面已经聊过模型转换、量化、算子融合但说实话那些优化手段最终都要落到硬件上执行。如果你不理解 GPU 怎么并行、显存数据怎么流动、线程怎么调度那很多优化对你来说就是个“碰运气”的开关——开了可能变快也可能崩掉。这篇文章就是这个系列的第 5 篇专门给部署工程师和算法工程师补一堂“GPU 底层课”。目标很明确让你理解 GPU 硬件架构里跟推理性能直接相关的部分掌握 CUDA 编程的基本逻辑并能自己定位一个算子跑得慢的瓶颈到底在哪。这篇文章不打算从“GPU 的历史”讲起也不堆概念。我会从实际部署中遇到的性能问题出发一层层拆开 CUDA 和 GPU 架构最后落到可操作的排查和调优手段上。无论是刚入门的新手还是已经在做推理加速的工程师应该都能从中找到能直接拿去用的东西。1. GPU 硬件架构抛开玄学直击推理性能的核心变量1.1 从 SM 到 CUDA Core算力是怎么组织起来的先纠正一个常见误区一张显卡并不是一堆核心简单地堆在一起。NVIDIA GPU 的基本计算单元叫SMStreaming Multiprocessor流式多处理器整张 GPU 是由几十甚至上百个 SM 组成的。每个 SM 内部又包含若干 CUDA Core也就是 FP32 运算单元、Tensor Core、专用缓存和调度器。举例来说A100 有 108 个 SM每个 SM 里有 64 个 FP32 CUDA Core 和 4 个第三代 Tensor Core。所以你在规格表上看到的“6912 个 CUDA Core”实际上是 108 × 64 算出来的。这个组织方式决定了并行度的上限要让 GPU 跑满你至少得让每个 SM 都有活儿干每个 CUDA Core 都有数据算。这里要特别提一下Tensor Core。从 Volta 架构开始加入的 Tensor Core是专门为矩阵乘加设计的硬件单元。它一次能算 4×4 甚至更大的矩阵运算吞吐量是普通 CUDA Core 的很多倍。模型推理里的卷积、全连接、Attention 本质上都是矩阵运算所以 Tensor Core 几乎决定了推理性能的天花板。这也是为什么混合精度推理FP16、BF16能带来巨大加速——FP16 的计算可以直接喂给 Tensor Core而 FP32 只能走普通 CUDA Core。搞部署的人不需要会写 Tensor Core 的汇编但要明白一件事算子能不能用到 Tensor Core取决于数据精度和内存布局。你调用的 cuBLAS、cuDNN 这些库会自动选择是否启用 Tensor Core但前提是你的输入张量是 FP16/BF16 且内存连续。如果你把模型强行跑在 FP32 下再好的显卡也发挥不出性能。1.2 显存带宽与内存层次推理模型卡在“搬数据”而不是“算数据”算力重要但还有一个更容易被忽略的指标——显存带宽。很多推理场景下瓶颈根本不是 GPU 算不快而是数据来不及送到计算单元里。打个比方。CPU 推理就像一个人去仓库搬货货架很远每一步都要走过去拿所以速度受限于走路时间。GPU 推理像一条传送带流水线大量工人CUDA Core站在流水线旁边传送带显存带宽源源不断把货送过来工人只需要伸手接住并计算。如果传送带太窄就算工人再多整体吞吐也上不去。现代显卡的显存带宽已经做到了几百 GB/s 甚至 TB/s 级别但这只是“峰值”。实际能达到多少取决于你的内存访问模式。如果线程访问的数据地址是连续的合并访存coalesced memory access显存能一次把一大块数据搬进来如果每个线程都去访问分散的地址带宽利用率可能连峰值的 10% 都不到。在 GPU 的存储体系里从快到慢大致是寄存器每个线程私有速度最快Shared Memory共享内存同一个 Block 内线程共享延迟只有几十个时钟周期L1/L2 缓存硬件自动管理命中率对性能影响很大全局内存显存容量大但延迟高大概几百个时钟周期这就是为什么很多算子优化都要用tiling分块策略先把数据从显存搬一块到共享内存然后线程们从共享内存里反复读取计算减少直接访问显存的次数。你去看 cuBLAS 的矩阵乘法实现核心思路就是这个。1.3 为什么显卡的“理论算力”不等于实际推理速度规格表上 A100 的 FP16 算力是 312 TFLOPS但你在实际推理中很少能跑到这个数字的一半。原因是多方面的第一kernel launch 开销是真实存在的。每个 CUDA kernel 从 CPU 发起要经过驱动层几微秒的开销对于动辄几毫秒的推理来说不算什么但如果你把一个大算子拆成了几百个小算子累计开销就会很可观。这也就是为什么算子融合比如把 ConvBNReLU 合成一个 kernel能带来显著加速。第二数据搬运往往比计算更耗时。推理时输入要从 CPU 内存拷贝到显存输出要拷回来如果模型结构里频繁出现小的数据传输这部分时间会占据很大比例。我把这个叫做“搬运税”你的实际帧率要扣掉这层税。第三occupancy占用率决定隐藏延迟的能力。GPU 通过切换线程束Warp来隐藏内存延迟。如果当前活跃的 Warp 数量太少当一个 Warp 在等待显存返回数据时SM 里没有其他 Warp 可供调度计算单元就会空转。占用率低算力再高也白搭。所以判断一个算子是否性能达标先看两个维度它是不是 memory-bound带宽瓶颈还是 compute-bound算力瓶颈。判断方法后面实操部分会细讲但记住一句话先分清瓶颈类型再决定优化手段。如果是带宽瓶颈单纯增加算力毫无意义你要做的是优化数据复用和访问模式如果是算力瓶颈才需要考虑混合精度、Tensor Core 或者算法层面的重写。2. CUDA 编程模型线程、内存与核函数的基本逻辑2.1 线程层次Grid、Block、Thread 和 Warp 的“三级组织”CUDA 的线程组织有点像公司层级整个公司Grid分成几个部门Block每个部门里有若干员工Thread。硬件调度时GPU 并不会一个一个 Thread 地执行而是以Warp通常 32 个线程为单位。一个 Block 会被拆成多个 Warp比如一个 Block 有 256 个线程那就是 8 个 Warp。理解这一层很关键因为它直接影响你写 kernel 时怎么安排线程、怎么设 Block 大小。经验值是这样的Block 大小一般取 128 到 512 之间别太大也别太小Block 的线程数最好取 32 的倍数因为一个 Warp 是 32 线程Grid 的大小也就是 Block 的数量应该足够多保证所有 SM 都“吃饱”从部署工程师的角度来看你不需要手写特别复杂的 kernel大部分时候调用 cuBLAS、cuDNN 就够了。但当你需要写 custom kernel比如融合 Attention、自定义激活函数时这些概念就是基本功绕不开。举个例子我做过一个 FlashAttention 类似的融合 kernel就是把 Q、K、V 的切分和 softmax 的在线计算合到一个 kernel 里。如果不懂 Block 和 Warp 怎么映射到矩阵分块根本无从下手。很多时候你从 GitHub 上抄一个 kernel 跑得慢就是因为没理解线程布局和数据划分的关系。2.2 内存访问与 shared memory把慢显存变成快缓存内存访问模式是 kernel 性能的第一杀手这一点怎么强调都不过分。最理想的情况是合并访存同一 Warp 内32 个线程访问的地址是连续的。这样 GPU 的显存控制器可以一次事务把 32 个数据全部取回。反过来如果 32 个线程访问的是完全分散的地址比如每次跳 8 字节那么一次事务只能拿到少量数据带宽浪费严重。还有一个容易忽略的点bank conflict共享内存冲突。共享内存被分成了 32 个 bank每个 bank 每时钟周期只能响应一个访问请求。如果多个线程同时访问同一个 bank 的不同地址就会发生冲突访问会被串行化性能骤降。最常见的就是跨步访问比如二维数组按行还是按列读取。说这些不是为了吓唬你而是为了让你明白很多“感觉不对劲”的性能问题本质是内存模式问题而不是计算问题。我在优化一个 GELULayerNorm 融合算子时最初版本运行时间 1.2ms改成按行连续读取并避免 bank conflict 后直接降到 0.3ms算子逻辑完全没变只改了内存访问方式。2.3 一个最简单的 CUDA 程序到底在干什么为了把上面这些概念串起来用一个最简单的 Vector Add 展示 CUDA 程序的完整流程。假设有两个长度为 N 的数组 a 和 b要计算 c a b。__global__ void vector_add(float *a, float *b, float *c, int n) { int idx blockIdx.x * blockDim.x threadIdx.x; if (idx n) { c[idx] a[idx] b[idx]; } } int main() { float *d_a, *d_b, *d_c; int n 1 20; size_t bytes n * sizeof(float); // 1. 在显存上分配空间 cudaMalloc(d_a, bytes); cudaMalloc(d_b, bytes); cudaMalloc(d_c, bytes); // 2. 把数据从 CPU 内存拷贝到显存 cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice); cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice); // 3. 启动 kernel int threads 256; int blocks (n threads - 1) / threads; vector_addblocks, threads(d_a, d_b, d_c, n); // 4. 把结果拷回 CPU 内存 cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost); // 5. 释放显存 cudaFree(d_a); cudaFree(d_b); cudaFree(d_c); return 0; }这段代码里有几个关键点值得注意第一线程索引的计算blockIdx.x * blockDim.x threadIdx.x是 CUDA 最基础的“套路”。它把一维 Grid 里所有线程映射到数组下标上。N 可能不是 blockDim.x 的整数倍所以要加一个if (idx n)判断否则会越界访问。第二cudaMemcpy的调用是同步的CPU 会阻塞等待数据拷贝完成。频繁地 Host 和 Device 之间拷贝数据是性能杀手。绝大多数推理优化都会尽量把数据一次性搬进显存减少来回传输。第三cudaMalloc开销很大不要频繁调用。在推理服务里通常的做法是预先分配好显存池运行时复用而不是每次请求都申请释放。这也是为什么 vLLM、TensorRT 等框架里都有显存池管理机制。3. 实操从环境检查到性能分析跑通第一个部署级 CUDA 流程3.1 环境自检驱动、Toolkit、PyTorch 版本怎么对齐部署中遇到的大部分 CUDA 问题根源都是环境版本不匹配而不是代码问题。我见过太多人折腾了一整晚最后发现是驱动和 Toolkit 版本对不上。先解决最基础的搞清楚驱动版本和 CUDA 版本的关系。运行nvidia-smi右上角会显示“CUDA Version: 12.x”这个数字表示当前驱动最高支持的 CUDA 版本。你再运行nvcc --version看到的才是你当前编译环境里的 CUDA Toolkit 版本。这两个版本可以不一样Toolkit 版本只要不高于驱动支持的版本就能正常工作。举个例子你的驱动支持的 CUDA 是 12.4你装的 Toolkit 是 11.8那编译 CUDA 代码没问题反过来你的驱动是 11.xToolkit 却是 12.x就一定会出错。我的建议是在部署环境里尽量保持“驱动支持版本 Toolkit 版本 深度学习框架版本”并且不要频繁升级驱动。驱动是操作系统的底层组件升级有风险而 CUDA Toolkit 可以通过 conda 或本地目录安装多份互相隔离。这里分享一个多版本 CUDA 共存的方法。CUDA Toolkit 安装时会提供安装到自定义目录的选项比如/usr/local/cuda-11.8和/usr/local/cuda-12.4。然后用软链接切换ln -s /usr/local/cuda-11.8 /usr/local/cuda export PATH/usr/local/cuda/bin:$PATH export LD_LIBRARY_PATH/usr/local/cuda/lib64:$LD_LIBRARY_PATH下次要切到 12.4只需把软链接改一下再刷新环境变量即可。这个方法比反复装卸 Toolkit 干净得多。很多开发机的显卡驱动只有一套但 CUDA 环境可以有好几个切换着用完全没问题。另外容器部署时最省心的做法是用官方 PyTorch 镜像镜像里已经配好了一整套 CUDA cuDNN 环境。你只需要保证宿主机的 NVIDIA 驱动有足够新的版本然后让容器把驱动暴露进去docker run --gpus all -it pytorch/pytorch:2.4.0-cuda12.1-cudnn9-runtime bash宿主机驱动的检测逻辑在 NVIDIA 容器工具包里它会检查容器的 CUDA 版本是否被宿主驱动支持。实际部署中我踩过坑镜像里是 CUDA 12.1但宿主驱动太老只支持到 11.x启动容器直接报 CUDA driver version is insufficient。升级宿主机驱动后问题消失。3.2 手写一个 Vector Add 并做性能对比这里做一个小实验帮助我们直观感受 GPU 程序的真实带宽和“为什么有时候 GPU 反而更慢”。N 取 1000 万长度分别用 CPU 单线程和 GPU kernel 各跑一次加法。CPU 版本大约耗时几十毫秒GPU 版本刨掉数据拷贝时间kernel 本身只要几十微秒。但把数据拷贝时间加上后你会发现整个 GPU 流程并没有快多少甚至在小数据量下可能更慢。原因就是刚才说的“搬运税”。数据传输的时间往往会主导整个流程这解释了为什么模型部署时一般不做逐层推理再拷回结果而是尽量把全模型放在 GPU 上全过程留在显存里。对纯计算型算子还有一个更敏感的问题第一次调用 CUDA kernel 会有上下文初始化和编译缓存所以实际测试性能时不能只跑一次要预热后取多次平均值。3.3 用 Nsight Compute 定位瓶颈读哪些指标手写 kernel 是一回事知道它为什么快、为什么慢是另一回事。在部署优化里用得最多的工具是Nsight Compute命令行是 ncu和Nsight Systemsnsys。先区分两个工具的定位nsys侧重整体时间线定位“时间花在了哪个阶段”kernel 还是 memcpyncu深入单个 kernel分析“SM 利用率、带宽利用、占用率”等微观指标实操时我会先跑 nsys 看整体时间分布如果发现某个 kernel 特别耗时再针对它跑 ncu 看细节。ncu 的典型用法ncu --set full ./my_program跑完后重点关注几个指标SM 利用率如果很低说明计算单元大量空闲多半是线程不足或访存等待Memory Throughput如果接近 100%说明算子已经打满显存带宽属于 memory-boundCompute Throughput如果很高而 Memory Throughput 低则是 compute-boundAchieved Occupancy实际占用率数值低说明 Warp 数量不足无法隐藏延迟这里有一个非常实用的判断方式如果 Memory Throughput 很高而 Compute Throughput 很低那么这个算子的优化方向就是减少显存访问比如用共享内存做数据复用。反过来如果 Compute Throughput 高则考虑混合精度、让计算走 Tensor Core。我在优化一个自定义的 RMSNorm kernel 时ncu 显示 Memory Throughput 大约是 92%说明瓶颈就是带宽。后来我改成了半精度输入带宽压力减半耗时几乎也减半验证了这个判断。4. 推理部署中典型的 CUDA 问题与排查实录4.1 CUDA out of memory 与显存碎片化“CUDA out of memory”可能是部署中最常见的报错但很多人不知道 OOM 其实分两种一种是真的显存不够另一种是显存碎片化导致无法分配连续大块内存。推理模型的显存大头一般是权重和 KV Cache。如果模型权重本身接近显存上限最简单的办法是量化从 FP16 切到 INT8 或 INT4显存直接减半或减到四分之一。如果 KV Cache 太大可以限制并发数或者用 PagedAttention 这种动态分配方案。碎片化的判断办法查看nvidia-smi里的显存占用发现空闲显存足够多但程序一跑还是 OOM多半是碎片化。原因是推理过程中不断有 tensor 被分配和释放显存空间被切成很多小块而 CUDA 申请连续大内存时找不到足够大的连续空间。解决办法也很粗暴推理时预先分配好显存池后续运算从池子里取不还给驱动。或者直接换用 TensorRT、vLLM 这类自带显存池的框架省得自己造轮子。4.2 “sm_120 not compatible” 类的算力不匹配很多人在新显卡上部署老模型时碰到过类似报错nvidia geforce rtx 5070 laptop gpu with cuda capability sm_120 is not compatible with the current cuda driver这个问题的本质是你的 CUDA Toolkit 版本太老不认识新一代显卡的计算能力编号compute capability。比如 Blackwell 架构的显卡计算能力是 12.0sm_120而旧 Toolkit 可能只支持到 sm_90 或更早。排查步骤很清晰运行nvidia-smi确认驱动是否支持这块显卡。如果驱动太老先升级驱动查看当前 CUDA Toolkit 版本nvcc --version到官方文档查这块显卡需要的 compute capability如果 Toolkit 版本过低安装新版本 Toolkit或者在编译时加-archsm_120参数注意PyTorch 这类框架是预编译的它自带了一段针对特定架构的 kernel。如果你的显卡架构比 PyTorch 预编译时支持的更新PyTorch 会尝试用 JIT 方式编译 kernel但前提是环境里得有匹配的 nvcc。所以遇到这种报错时最稳的解法是升级 PyTorch 到支持新架构的版本并把它对应的 CUDA 版本装上。4.3 驱动、容器、WSL2 里的常见坑容器化部署已经成为主流但容器里的 GPU 问题也有不少坑这里提几个高频问题第一宿主机装了驱动容器里却看不到 GPU。这种情况多半是 NVIDIA Container Toolkit 没装好。注意NVIDIA 驱动只在宿主机上装即可容器里不需要装驱动但需要暴露设备。检查/etc/docker/daemon.json里有没有配置好 nvidia runtime以及是否用了--gpus all启动。第二WSL2 里装 CUDA。很多开发机用 Windows WSL2 跑 GPU 任务踩坑主要在驱动上。WSL2 里的 GPU 驱动是装在 Windows 侧的WSL 内部不要也不能再装驱动只需要安装与驱动匹配的 CUDA Toolkit 即可。如果你在 WSL 里执行nvidia-smi能看到显卡说明驱动通路已经正常。第三多用户共享 GPU 时权限问题。nvidia-smi能显示设备但用户运行时提示找不到 GPU大部分是/dev/nvidia*设备文件的权限问题把用户加入 video 组一般能解决。4.4 常见 CUDA 部署错误速查表报错信息根本原因快速排查手段CUDA out of memory显存真正不足或碎片化nvidia-smi 查看空闲显存分布减少 batch 或换量化CUDA driver version is insufficient驱动的 CUDA 支持版本低于运行时需求nvidia-smi 查看右上角支持版本升级驱动sm_120 is not compatibleToolkit 或框架版本过老不识别新架构升级 Toolkit、PyTorch编译时指定 -archlibcudnn.so.9: cannot open shared object filecuDNN 版本不对或未安装安装匹配版本的 cuDNN或用官方 Docker 镜像no kernel image is available for execution on the device编译时架构不匹配常见于自定义 kernel检查 -arch 参数交叉编译时加上目标架构WSL 里 nvidia-smi 报错Windows 侧驱动未装或版本过旧在 Windows 安装 NVIDIA 驱动不要在 WSL 内装驱动这张表更像是排障清单日常部署遇到问题先对着表过一遍能省不少时间。还有一个经验值得分享报错信息里往往藏着版本号先看版本再看代码。每次排查问题第一步永远是收集环境信息驱动版本、Toolkit 版本、cuDNN 版本、PyTorch 版本、显卡型号。把这几项列出来对照官方兼容表80% 的问题都能直接定位。结语我个人在做模型部署这几年的体会是环境问题比代码问题多带宽问题比算力问题多架构理解问题比工具问题多。很多人在 GPU 上跑模型就完事了不去深究为什么快为什么慢直到遇到性能指标上不去才回头补课。如果你现在正处在“优化不知道怎么下手”的阶段我的建议是别急着换模型换框架先用 ncu 跑一遍性能分析看清楚瓶颈在哪个方向再决定下一步动作。最后再分享一个小技巧在你常用的部署环境里做一个print_env.sh脚本把显卡、驱动、CUDA、cuDNN、PyTorch、容器的信息一次性打印出来。排查问题、写 issue、跟同事沟通时这个脚本能帮你节约大量时间。GPU 部署这条路没有捷径但把基础打扎实之后后面每一步的回报都是实打实的。