1. 为什么做模型部署的人必须懂CUDA和GPU架构1.1 推理延迟和吞吐量背后到底是谁在干活做模型部署这几年我最大的感受是很多人把模型部署理解成“把训练好的模型加载起来调一下batch size开个服务等请求”但一旦遇到真实的性能问题比如首token延迟压不下去、GPU利用率上不去、并发一高就超时就完全不知道从哪里下手了。这些问题的核心最终都会落到底层硬件身上。你用的是GPUGPU的指令怎么执行、数据怎么搬运、线程怎么调度这些全由CUDA编程模型和GPU架构共同决定。哪怕你用的是PyTorch、TensorRT、ONNX Runtime这种高层工具它们在下层编译和优化时无一例外都在做同一类事情把计算组织成适合GPU执行的形态。你要是完全不懂这套底层逻辑就只能靠试参数碰运气。举个我实际遇到的例子。之前调一个BERT类模型的推理服务单条请求的延迟一直卡在12毫秒左右怎么压都下不去。我用Nsight Systems抓了一下时间线发现模型里有个自注意力层矩阵乘加softmax这一步在GPU上的执行时间居然占了将近一半而大多数SM流式多处理器的占用率只有30%不到。问题出在哪里算子粒度太小矩阵规模不够大启动开销和数据搬运耗时占比太高了。当时我用了最简单粗暴的办法——把相邻的几个小算子手动做算子融合合并成一个kernel延迟直接从12毫秒掉到了8毫秒以下。这件事让我意识到高层框架给不了你的恰恰是这些最底层的性能洞察。1.2 不懂CUDA能调好推理性能吗先说结论如果你只是常规调用现成推理引擎不懂CUDA也能勉强把活干完但一旦涉及性能调优、内存优化、自定义算子、异构计算不懂CUDA寸步难行。我自己带过不少新人遇到过一个挺典型的场景有人写了个预处理逻辑把图片缩放和归一化放在CPU上用OpenCV做再把处理完的张量传到GPU上500毫秒的推理延迟里有300毫秒花在了数据传输和CPU同步上。他当时的直觉是“CPU反正闲着也是闲着能分担一点是一点”但完全没意识到CPU和GPU之间存在PCIe总线带宽瓶颈来回传输一次数据的时间可能比GPU本身跑算子还贵。这种问题光靠调框架配置是解决不了的。你要知道GPU的架构特点比如SM的调度方式、内存带宽和计算吞吐之间的比例关系、PCIe传输的固定开销才知道为什么“看起来合理的方案”实际上是在帮倒忙。所以这篇内容我打算从一个经常跑模型部署和推理优化的人的角度把CUDA编程模型和GPU架构那些最基础但又最核心的东西掰开讲清楚。所有概念我都会用实际推理场景里的例子来讲不会只停留在教科书式的定义上。读完你至少能搞清楚GPU到底是靠什么把速度拉上来的、CUDA的线程模型是怎么运作的、为什么内存访问比计算更容易拖垮性能以及在实际部署中应该按什么思路去调优。2. GPU架构基础从“疯狂的多核加速器”说起2.1 从黑盒到半透明你需要知道的GPU内部结构把GPU想成一个大型超市配送中心。CPU是一个全能型选手什么活都能干但手脚慢一次只能同时处理几件事。GPU不是GPU是一群数量庞大但能力单一的搬运工每一个都不算聪明但架不住人多而且干同一种活的时候配合极其默契。从物理结构上说一块NVIDIA GPU由若干个**流式多处理器Streaming MultiprocessorSM**组成。每个SM内部有一批CUDA核心专业点叫“算术逻辑单元”负责浮点计算以及共享内存、寄存器文件、调度器等资源。以常见的A100为例108个SM每个SM有64个FP32 CUDA核心总共6912个核心。RTX 3090则是82个SM每个SM 128个FP32核心总共10496个核心。注意这个“核心”和CPU那个核心不是一个概念CUDA核心更接近CPU里的ALU它的职责就是执行单个浮点运算或整数运算本身没有完整的指令控制能力。关键点在于SM是GPU执行计算的基本单位。你把一个任务丢给GPU驱动和CUDA运行时会把这个任务拆成若干个小块分配到不同的SM上。每个SM内部再把这一个小块进一步拆成线程束warp来调度执行。线程束是GPU调度的真正粒度——32个线程为一组SIMT单指令多线程模式下这32个线程同时执行同一条指令只是各自处理不同的数据。这个设计初衷就是为了让大量核心在同一时刻保持“整齐划一”的忙碌状态避免出现有些核心等指令、有些核心空闲的情况。这套硬件设计对模型推理有几个直接影响。第一矩阵乘法类算子天然适合GPU矩阵乘法本质是大量独立的乘加运算可以完美地映射到大量并行线程上。第二分支逻辑越少越好同一个线程束里只要有线程走不同的分支这个线程束就会被迫串行执行所有分支路径性能会受到明显影响。所以你在写CUDA kernel时要尽量避免类似if (threadIdx.x % 2 0) {...} else {...}这种按线程ID分叉的逻辑。2.2 从Ampere到Hopper再到Blackwell架构演进对推理的影响选GPU硬件时很多人只看显存大小和价位很少关注架构代际差异。但架构差异对推理性能的影响远比你想象的大。以最近几年最常用的三代架构来说架构代表显卡关键改进对推理优化的直接影响AmpereA100, RTX 30系引入第三代Tensor Core支持TF32和稀疏化FP16矩阵运算吞吐大幅提升TF32可以在不损失太多精度的情况下跑FP32精度的网络HopperH100, RTX 40系第四代Tensor CoreTransformer引擎DPX指令针对Transformer结构做了指令级优化FP8支持让推理显存和带宽压力骤降BlackwellB200, RTX 50系第五代Tensor Core第二代Transformer引擎2倍于Hopper的FP4推理吞吐但兼容性问题也会更突出新一代架构通常会引入新的计算精度格式如TF32、FP16、BF16、FP8、INT8这些格式直接决定了推理时的吞吐上限。比如你用FP16代替FP32做推理A100上的FP16 Tensor Core吞吐是FP32的2倍H100的FP8吞吐又是FP16的2倍。这不是简单的“精度砍半、速度翻倍”而是硬件上Tensor Core本身就是为低精度高并行设计的。所以模型部署里的精度量化本质上是赌硬件的低精度计算单元能带来多大的加速收益。这里有一个很多人忽略的点量化带来的加速收益更多来自内存带宽节省而非计算加速。尤其在Decoder推理自回归生成场景下模型是batch size很小的逐token生成计算强度很低瓶颈几乎全在显存带宽上。把权重从FP32砍到FP8意味着同样时间内能从显存里读出4倍的数据这才是量化在LLM推理中收益巨大的根本原因。2.3 Tensor Core推理优化的关键加速单元Tensor Core是NVIDIA从Volta架构开始引入的专用矩阵乘单元它本质上就是一个“一次算一个4x4或1x4矩阵乘”的专门硬件块。普通CUDA核心算矩阵乘是一个个乘加指令循环执行Tensor Core则是直接把一个小矩阵整体吞进去一次输出一整块矩阵乘结果。为什么这很重要因为深度学习推理里绝大多数计算量都集中在卷积和矩阵乘上。以Transformer模型为例QKV投影、注意力分数计算、输出投影全都是矩阵乘。Tensor Core的出现相当于给这些算子修了一条高速公路。在推理优化时Tensor Core的利用情况是必须盯住的指标。用Nsight Compute看kernel的分析报告时如果你看到Tensor Pipe UtilizationTensor Core利用率只有百分之十几说明kernel没有吃的Tensor Core红利该检查数据排布和分块策略了。通常优化思路有这么几层确保数据精度匹配Tensor Core要求A100支持TF32、FP16、INT8H100/B200额外支持BF16、FP8、FP4尽量用官方库cuBLAS、cuDNN而不是自己写通用矩阵乘自己写kernel时分块大小要对齐Tensor Core的输入尺寸要求通常都是用16x16或16x8的分块我在实际部署中见过一个很典型的案例同一个矩阵乘算子用cuBLAS跑只要0.3毫秒用PyTorch默认的优化版约0.5毫秒而如果有人用基础CUDA kernel没做tiling直接裸写耗时可能直接飙到5毫秒以上。差距有10倍以上原因就是没有利用好Tensor Core和共享内存。3. CUDA编程模型软件怎么编排硬件干活3.1 网格、线程块、线程三级结构到底怎么理解CUDA编程模型的核心是一个三级线程结构网格Grid→ 线程块Block→ 线程Thread。对应到GPU硬件上一个Grid里的所有线程块会被分发到不同的SM上执行一个线程块里的所有线程会在同一个SM内部执行。线程块是软件上最小编排单位SM资源共享内存、寄存器、线程数上限决定了一个SM最多能同时驻留多少个线程块。这里有个经常被误解的概念线程块内部才能用共享内存和block级同步不同线程块之间没法直接同步也不该互相通信。如果你把一个需要跨块协作的算法硬塞进来要么需求本身就是不合理的应该重新设计数据切分要么需要额外用全局内存加原子操作或网格级同步CUDA 9之后有cooperative groups支持但性能代价很大。实际工作中我判断线程组织结构通常就三个标准让每个线程负责的工作量尽量平均让线程块大小是32的倍数32线程1个线程束避免线程束内出现空槽位总线程数不要远大于实际需要的并行工作量避免过度创建线程导致调度开销比如你要在GPU上一个长度100万的向量做逐元素操作最省事的做法是开256个线程的线程块总共3907个线程块。每个线程处理一个元素用gridDim.x * blockDim.x threadIdx.x算出自己的全局索引。块的个数略多于SM的驻留能力没关系GPU的块调度器会自动分配但太多的小块反而增加调度开销。之前调一个向量Add kernel把块大小从128调到256耗时掉了差不多15%原因是块太少时每个SM隐藏访存延迟的能力变弱了。3.2 线程束是性能调优的基本观察单位线程束这个概念我觉得是CUDA性能最核心的东西没有之一。GPU的世界里SM每次都是以线程束为单位取指令、派发指令的。一个线程块里有256个线程SM实际会把它拆成8个线程束每32个线程一个束来调度。硬件上有一个隐藏延迟的机制叫线程束切换。当一个线程束执行内存加载这种高延迟操作时SM会立刻切换到其他准备好的线程束继续执行计算指令等内存数据回来了再切回来。所以对于GPU来说足够的并行度 隐藏延迟的资本。这也解释了为什么GPU跑深度学习模型时batch size越大单样本吞吐越高。batch大→每层算子的矩阵乘规模大→并行度足→SM能通过线程束切换把访存延迟藏得干干净净。反过来batch为1的LLM解码阶段矩阵乘规模其实很小比如一个1x512的向量乘512x512的权重矩阵并行度不足SM的调度器经常空转等待显存返回数据。一种比较有效的优化方案是动态批处理dynamic batching。把同一时刻到达的多个请求攒在一起凑成一个更大的矩阵乘一次喂给GPU。这在LLM推理服务里已经是事实标准了。vLLM的continuous batching就是这么干的。理解了线程束调度原理你就能明白为什么这个设计能带来巨大收益而不是仅仅“看起来更合理”。3.3 写一个最简单的CUDA Kernel需要什么光看概念不够怎么上手写和编译CUDA代码我经过反复实践后认为最快的路线是装好CUDA Toolkit后用nvcc编译一个向量加法程序。__global__ void vecAddKernel(const float* a, const float* b, float* c, int n) { int idx blockIdx.x * blockDim.x threadIdx.x; if (idx n) { c[idx] a[idx] b[idx]; } }编译nvcc -O3 -archsm_80 -o vec_add vec_add.cu这里-archsm_80表示针对Ampere架构A100/RTX 30系优化。如果你是H100或RTX 40系用sm_90B200或RTX 50系用sm_120。不过要注意sm_120目前和很多老库存在兼容性问题如果编译时提示“cuda capable sm_120 is not compatible”要么降级到sm_100要么检查PyTorch/TensorRT的CUDA支持版本。我自己在RTX 5070 Laptop上踩过这个坑最新PyTorch也要到2.6以上才支持Blackwell架构。编译完成后用cudaMalloc分配显存cudaMemcpy把数据从CPU拷贝到GPU执行kernel再拷贝回CPU。这基本是早期学习时最常见的“程序人生显存检查、流管理、内存拷贝”。CUDA C对运行时检查很严格大部分已知的CUDA error都要靠查返回值和处理cudaError_t类型去定位少了这步代码出错时经常会看到随机性的数据错乱很难排查。更贴合实际部署的做法是直接用PyTorch的torch.cudaAPI写自定义kernel借助它的CUDA封装简化掉繁琐的设备内存管理流程。一个PyTorch自定义CUDA算子通常只写一个CUDA C扩展用pybind11和torch.utils.cpp_extension.load_inline机制加载。对刚上手的人来说这个路线比纯C开发简单不少又比纯Python循环快几个数量级。我强烈推荐先用这种方法跑通一个完整流程再回头啃核心里更深入的优化技巧。4. 实操从零优化一个CUDA算子手把手调一遍4.1 环境准备和验证别一上来就踩版本坑CUDA开发和模型部署最大的坑是环境不匹配这句话我必须放在最前面说。CUDA Toolkit、NVIDIA驱动、PyTorch/TensorRT这三者各自对版本有严格约束。日常接到报错私信里超过一半是版本搭配导致的。比较省心的搭配经验是这样先装驱动驱动版本决定你能支持哪些CUDA的最高版本再装CUDA Toolkit可以同时装多个版本习惯上用sudo apt install cuda-toolkit-12-4这种装12.412-8装12.8切换用update-alternatives或设置CUDA_HOME环境变量然后装PyTorch用官方pip install torch --index-url https://download.pytorch.org/whl/cu124这种方式保证PyTorch里的CUDA版本匹配你的Toolkit最后确认一切OKpython -c import torch; print(torch.version.cuda, torch.cuda.is_available())验证GPU驱动的经典命名是nvidia-smi不是nvidia-smi -a。看右上角的CUDA Version它代表当前驱动能兼容的最高CUDA运行时版本并不代表你机器上实际装了那个版本的Toolkit。很多人以为这个数字就是当前CUDA环境版本结果版本对应不上一堆稀奇古怪的报错。如果你要用容器形态部署模型直接在Docker镜像里装CUDA全量Toolkit很浪费镜像体积建议直接用官方的nvidia/cuda:12.4.1-runtime-ubuntu22.04镜像只需要一个运行时就够了里面包含了运行CUDA程序所需的所有动态库libcudart.so、libcublas.so等编译工具链和头文件不包含在内。开发编译阶段才需要devel镜像。4.2 一个具体案例优化向量归约Reduction算子向量归约把一个大数组求和成一个值是很多推理场景的基础算子像LayerNorm、Softmax里的求和、attention分数归一化都离不开它。这个算子看起来简单要调好其实不容易是个特别好的练手对象。基础版本写起来很简单__global__ void sumKernel(const float* input, float* output, int n) { int idx blockIdx.x * blockDim.x threadIdx.x; double sum 0.0; for (int i idx; i n; i gridDim.x * blockDim.x) { sum input[i]; } // 用原子加把当前线程的部分和累加到output if (threadIdx.x 0) { atomicAdd(output, (float)sum); } }这个版本有多慢实测下来长度为1000万的数组求和跑出来大约需要0.8毫秒。听起来不多但放进LayerNorm里网络每一层都要跑好几遍这种操作累积下来对推理延迟影响不小。优化的第一步是树形归约。每个线程块内部先把线程的值规约成一个数再用原子操作做块间合并而不是万万个线程全去抢同一个地址做原子加。原因是原子操作本质上是在做串行化大量线程同时往同一个内存地址累加性能会被锁冲突拖垮。__global__ void sumKernelTree(const float* input, float* output, int n) { __shared__ float shared[256]; int idx blockIdx.x * blockDim.x threadIdx.x; float val (idx n) ? input[idx] : 0.0f; shared[threadIdx.x] val; __syncthreads(); // 树形归约stride思想 for (int stride blockDim.x / 2; stride 0; stride 1) { if (threadIdx.x stride) { shared[threadIdx.x] shared[threadIdx.x stride]; } __syncthreads(); } if (threadIdx.x 0) { atomicAdd(output, shared[0]); } }这里有个很容易被新手忽略的点__syncthreads()必须放在每个线程都执行到的地方否则线程块会死锁。你可以想象成全班一起做操一个人不做完当前动作其他人就得一直等他那个不做动作的同学如果在后面直接跳到了下一轮前面的人就永远等不齐了。树形归约版本实测在同样数据量下可以跑到0.2毫秒左右是朴素版本的4倍。但还可以继续优化——最后一个线程块做最终归约完全省掉原子操作。做法是每个线程块归约完自己的部分和之后把结果写到一个全局数组的概率比如blockDim.x个位置然后让一个单独的线程块接管最终合并。这样原子操作一次都不需要发生性能还能再提15%~20%。这个优化过程能说明一个非常关键的观点写CUDA kernel时性能瓶颈往往是并发控制原子操作、同步而非算术本身。很多人在写kernel时只顾着算数却把所有并发问题都丢给了硬件解决结果性能上不去。4.3 用Nsight工具定位瓶颈而不是靠猜我在做推理优化时最常被问到“你是怎么知道该优化哪里”答案很简单用profiler跑一遍让数据告诉我不是靠猜。NVIDIA官方提供了三件套工具nvidia-smi看实时GPU利用率和显存占用适合做整体观察Nsight Systems全系统级别的性能分析能看到CPU和GPU的时间线、kernel启动间隙、数据传输时间适合定位宏观问题比如CPU预处理太慢导致GPU空转Nsight Compute单kernel级别的微架构分析能看到每个kernel的SM利用率、内存吞吐、线程束状态适合精调kernel内部实现Nsight Compute报告里有一个“Speed of Light”速度之光指标图会显示SM计算单元吞吐和内存吞吐的百分比。如果SM throughput用满了而Memory throughput很低说明计算是瓶颈反过来则说明访存是瓶颈。还有的kernel两边都用不满那大概率是并行度不足或存在同步等待问题。用这个工具看归约kernel你会发现基础版本的主要瓶颈在原子操作串行化而树形归约版本主要瓶颈就转移到了最终一次全局写和线程束分歧上。这时候再用__shfl_down_sync做线程束内归约warp shuffle可以做到一个线程束内通过寄存器直接交换数据完全绕过共享内存进一步提升归约速度。这个优化到位后向量归约的耗时能从0.2毫秒再压到0.1毫秒左右。对你没有看错同一份逻辑从粗放到精细优化能差8倍以上的性能。5. 内存层级与访存优化推理性能的真正胜负手5.1 GPU内存金字塔到底哪一层负责什么很多开发者刚接触GPU时觉得显存就是显存直到第一次发现kernel跑得慢去看profiler才发现问题全出在访存身上。GPU内部的内存层级其实是分得非常明确的每一级的速度和容量差别极大。从快到慢粗略排一下存储层级容量相对速度A100示例作用寄存器文件每个SM 256KB最快几乎无延迟每个线程私有的最高速存储共享内存每个SM 192KB几十TB/s级带宽线程块内共享速度接近寄存器L1/L2缓存L1 256KB/SML2 40MBL2约6~8TB/s自动缓存全局内存数据全局显存HBM数十GB约2TB/s主存储区所有线程可见主机内存数十GB~数百GBPCIe 4.0约32GB/s数据进出GPU的通道这个速度差异比是什么概念全局内存的延迟大概400~600个时钟周期而寄存器访问只需要几个时钟周期差了差不多两个数量级。共享内存虽然快但容量极小一整个SM才几百KB。所以CUDA优化的本质就是在这么狭窄的高速存储空间里想方设法把最常用的数据塞进去把不常用的数据按块搬到全局内存再让所有线程尽量复用。我用一个生活化的例子解释GPU的全局内存像是一个大图书馆书都在馆里但每次去借书都要走一段很长的路几百个周期。共享内存像是研究室的阅览桌你可以把一批常用书搬到桌上几个组员围着这张桌子快速翻阅省去了反复跑图书馆的时间。5.2 共享内存Bank Conflict一个让新手一脸懵的坑共享内存虽然是“高速缓存”但它内部又分成32个Bank有32个存储体每次访问周期内32个线程即一个线程束可以同时访问共享内存。理想情况下如果每个线程访问的地址落在不同Bank里这个访问周期就能在一个周期内完成。但如果有两个线程同时访问同一个Bank的不同地址硬件就必须分两次完成——“Bank Conflict”冲突发生性能直接打对折。再具体一点如果32个线程同时访问同Bank的同地址比如都读共享数组的一个元素硬件是支持“广播”的不产生冲突。但如果16个线程访问Bank 0的地址A另16个线程访问Bank 0的地址B硬件就需要分两个周期。代码看起来没毛病性能却莫名衰减一半这种问题不靠Nsight Compute的Bank Conflict指标很难直观感觉到。处理Bank Conflict的经典技巧是填充padding。比如要把一个float data[32][32]的二维数组以行优先方式存储访问时每列刚好落在同一个Bank上就会导致Bank Conflict。改成float data[32][33]每行多一个float行的起始偏移就会错开Bank分布被打散冲突就解决了。这里的本质是通过加1个元素的偏移让本来整齐排列的存储地址在Bank上也能均匀落在不同的位置上。在模型推理里共享内存用得最多的场景就是矩阵乘法的分块优化。把大矩阵切分成小块比如16x16的小矩阵每个线程块负责一块先把数据从全局内存搬到共享内存然后每块数据被反复使用多次。A100上的经典优化中一个16x16的Tile会让每个浮点数被读取大约16次这极大减少了全局内存带宽的压力。这也就是为什么FP16、INT8精度能让推理变快的一个重要原因——数据量少一半搬数据的压力也跟着少一半共享内存能存下更多批次的数据。5.3 融合和访存模式优化部署里的实际收益来源模型部署优化中最落地的两个方向就是算子融合和访存方式优化。算子融合的出发点很朴素GPU里每个kernel的启动和全局内存读写都有固定开销两个相邻算子如果都写中间结果到全局内存、下一个再读整个全局内存等于白白多走两趟PCIe/HBM。把它们融合成一个kernel中间量直接留在寄存器或共享内存里省掉的访存时间非常可观。举一个具体例子在做Transformer推理优化时QK^T算完之后跟softmaxsoftmax又要做一次row-wise max和sum传统做法会有两个中间矩阵写入全局内存再分别读取。如果把这些全部融合到QK^T的kernel尾部直接用寄存器里的局部矩阵做softmax最后只输出P矩阵一次访存量直接砍掉一半。在大batch场景下这个融合能让单层Transformer的延迟降低20%以上。访存模式优化的另一面是合并内存访问coalesced memory access。GPU访问全局内存时一个线程束的32个线程最好访问一段连续的内存地址这样硬件可以把32个4字节数据组合成一笔较大批量的传输完成。如果线程的访问地址是打散的比如按列读一个按行存储的二维数组那么一个内存请求就会退化成多次小传输带宽利用率惨不忍睹。这个在高性能计算里叫“访存合并性”。在写kernel时最典型的提示是让线程ID和连续索引一一对应。比如数据是一个NxK的矩阵行优先存储要按行做处理时应该让线程ID对应列方向连续位置而不是让线程ID取不同行。很多新手一上来把二维网格的threadIdx.x直接当成行索引结果全局内存访问完全不连续性能可能差出十倍。6. 模型部署中的CUDA实战问题与排查经验6.1 常见问题速查表见了别慌这些年在模型部署和CUDA环境上踩过的坑太多这里整理一个高频问题速查表可以直接当作排错手册。现象大概率原因解决办法CUDA error: no kernel image is available for execution on the device编译时的-arch架构和实际GPU不匹配用-archnative或显式指定对应的sm_XX重编比如A100是sm_80RTX 4090是sm_89PyTorch报CUDA out of memory但nvidia-smi显存还有剩余PyTorch缓存分配器把显存预先占用了nvidia-smi看到的是已分配量不是实际使用量用torch.cuda.empty_cache()或设置PYTORCH_CUDA_ALLOC_CONFmax_split_size_mb:128DataLoader多进程导致CUDA显存翻倍每个workser进程都会初始化CUDA context显存被重复占用减少num_workers或者把数据预处理改成主进程内一次完成程序运行一段时间后突然报CUDA Illegal memory access数组越界、野指针或异步执行时的资源竞争在cudaMemcpy和kernel调用后加cudaGetLastError()检查用Compute Sanitizer复盘定位Windows下PyTorch与CUDA版本不匹配安装时用了默认cpu版本或者PyTorch版本太旧torch.__version__确认配cu118/cu124等带CUDA后缀版本同一个模型TensorRT比PyTorch快不了多少算子融合没做透或模型里有大量动态shape导致规划退化检查Builder的优化策略设置尽量固定batch size开启FP16/INT86.2 我踩过最深的三个坑顺便聊聊教训第一个坑是异步执行带来的数据竞争。CUDA里kernel执行和CPU端代码是异步的CPU提交kernel后立刻返回不会等GPU执行完。这看起来很好但如果你在kernel执行过程中去读它上面写的输出缓冲读到的就是脏数据。特别坑的是这种错误不一定是稳定复现的有时候跑10次出现1次排查起来如大海捞针。养成好习惯需要跨设备同步时用cudaDeviceSynchronize()或cudaMemcpy这个本身就是同步的不要依赖于“上一次跑成功了所以这次也会成功”。第二个坑是动态batch下显存碎片化。LLM推理服务经常要动态批处理每次请求大小不一致显存分配模式很乱。运行一天后明明总显存占用不高但申请一批连续大内存时就会失败。一开始我用的是PyTorch默认的分配器后来换了vLLM或TensorRT-LLM这类自带显存管理方案的框架或者自己实现一个简单的显存池问题就基本消失了。对自研服务来说预先分配一块大显存、以page方式管理分配比频繁调用cudaMalloc高效得多。第三个坑比较冷门但遇到的人不少Multi-Process ServiceMPS和并发kernel之间的资源挤压。我在调多路推理并发时为了提升GPU利用率开了MPS结果反而出现了kernel互相干扰、个别请求延迟飙升的情况。后来发现MPS的共享上下文对显存资源分配颗粒度影响不小并发量小的时候效果不明显并发量一旦上去反而会劣化。如果你的推理负载本身就是多路并发独立请求开MPS之前一定要先用关闭状态下的性能基线做对比。6.3 CUDA进阶还有哪些东西在推理场景里值得学当你把基础的线程组织、内存管理、算子融合都吃透了会发现在推理优化这条路上还有几块内容是联系在一起的。CUDA Graphs是我认为最被低估的技术之一。普通方式调用kernel时每个kernel的启动开销大约5~10微秒如果一个推理模型有100个算子光启动开销就0.5~1毫秒。在追求低延迟的场景这个开销占比很可观。CUDA Graphs的作用是把整个模型的计算图提前捕获提交一次就把一串kernel按依赖关系全部发射出去单次推理的启动开销直接被压缩到一个很低的水平。TensorRT里有专门的enqueueV3接口支持图模式PyTorch 2.x也支持torch.cuda.graphs拿来做推理延迟优化效果非常直接。**持久化kernelPersistent Kernel**则是另一种思路。常规kernel执行完就释放资源下一算子又要重新分配和调度。在推理流水线里可以让一个kernel一直驻留GPU循环处理多个请求减少重复的初始化开销。这种模式在正式的推理引擎里用得很多缺点是编程复杂度高、调优难度大建议先从CUDA Graphs入手打牢基础后再进阶。最后一定要提的是精度与数值稳定性问题。位码优化调得太猛显存压力和速度是上去了但模型精度掉了一个多点用户立刻感知到回答质量问题这在实际业务里属于重大事故。在推理部署时所有的精度压缩都要经过离线评测集验证并且要定义好“误差阈值”别只看损失函数直接拿线上真实数据分布来测试。写在最后的一点体会从最开始在PyTorch里直接掉kernel到后来自己维护一个推理引擎中间被无数隐晦的性能问题折磨过。现在回头看CUDA和GPU架构知识真正的价值不是让你丢掉PyTorch自己从头写所有算子而是当框架自带的优化手段不够用时你能清楚地知道瓶颈在哪个环节、用什么手段去解。我最常推荐给团队同事的思路是先用profiler把最耗时的三个算子列出来只看不猜然后判断瓶颈是访存还是计算看Speed of Light最后针对性地做算子融合、分块优化或精度调整。这套方法对新手和进阶者都适用本质上是一种“用数据找真相”的工作方式。如果你正准备把模型部署和推理优化当长期方向CUDA这门基本功值得花一年半载慢慢啃不需要背常用API但要把线程组织、内存层级、调度模型这几块吃透。地基不牢的优化基本靠运气地基牢了之后你会在心里慢慢形成一张“性能地图”拿到什么模型都能快速定位瓶颈所在。这是没有捷径、但足够扎实的路子。
模型部署必知：CUDA与GPU架构的性能优化核心
1. 为什么做模型部署的人必须懂CUDA和GPU架构1.1 推理延迟和吞吐量背后到底是谁在干活做模型部署这几年我最大的感受是很多人把模型部署理解成“把训练好的模型加载起来调一下batch size开个服务等请求”但一旦遇到真实的性能问题比如首token延迟压不下去、GPU利用率上不去、并发一高就超时就完全不知道从哪里下手了。这些问题的核心最终都会落到底层硬件身上。你用的是GPUGPU的指令怎么执行、数据怎么搬运、线程怎么调度这些全由CUDA编程模型和GPU架构共同决定。哪怕你用的是PyTorch、TensorRT、ONNX Runtime这种高层工具它们在下层编译和优化时无一例外都在做同一类事情把计算组织成适合GPU执行的形态。你要是完全不懂这套底层逻辑就只能靠试参数碰运气。举个我实际遇到的例子。之前调一个BERT类模型的推理服务单条请求的延迟一直卡在12毫秒左右怎么压都下不去。我用Nsight Systems抓了一下时间线发现模型里有个自注意力层矩阵乘加softmax这一步在GPU上的执行时间居然占了将近一半而大多数SM流式多处理器的占用率只有30%不到。问题出在哪里算子粒度太小矩阵规模不够大启动开销和数据搬运耗时占比太高了。当时我用了最简单粗暴的办法——把相邻的几个小算子手动做算子融合合并成一个kernel延迟直接从12毫秒掉到了8毫秒以下。这件事让我意识到高层框架给不了你的恰恰是这些最底层的性能洞察。1.2 不懂CUDA能调好推理性能吗先说结论如果你只是常规调用现成推理引擎不懂CUDA也能勉强把活干完但一旦涉及性能调优、内存优化、自定义算子、异构计算不懂CUDA寸步难行。我自己带过不少新人遇到过一个挺典型的场景有人写了个预处理逻辑把图片缩放和归一化放在CPU上用OpenCV做再把处理完的张量传到GPU上500毫秒的推理延迟里有300毫秒花在了数据传输和CPU同步上。他当时的直觉是“CPU反正闲着也是闲着能分担一点是一点”但完全没意识到CPU和GPU之间存在PCIe总线带宽瓶颈来回传输一次数据的时间可能比GPU本身跑算子还贵。这种问题光靠调框架配置是解决不了的。你要知道GPU的架构特点比如SM的调度方式、内存带宽和计算吞吐之间的比例关系、PCIe传输的固定开销才知道为什么“看起来合理的方案”实际上是在帮倒忙。所以这篇内容我打算从一个经常跑模型部署和推理优化的人的角度把CUDA编程模型和GPU架构那些最基础但又最核心的东西掰开讲清楚。所有概念我都会用实际推理场景里的例子来讲不会只停留在教科书式的定义上。读完你至少能搞清楚GPU到底是靠什么把速度拉上来的、CUDA的线程模型是怎么运作的、为什么内存访问比计算更容易拖垮性能以及在实际部署中应该按什么思路去调优。2. GPU架构基础从“疯狂的多核加速器”说起2.1 从黑盒到半透明你需要知道的GPU内部结构把GPU想成一个大型超市配送中心。CPU是一个全能型选手什么活都能干但手脚慢一次只能同时处理几件事。GPU不是GPU是一群数量庞大但能力单一的搬运工每一个都不算聪明但架不住人多而且干同一种活的时候配合极其默契。从物理结构上说一块NVIDIA GPU由若干个**流式多处理器Streaming MultiprocessorSM**组成。每个SM内部有一批CUDA核心专业点叫“算术逻辑单元”负责浮点计算以及共享内存、寄存器文件、调度器等资源。以常见的A100为例108个SM每个SM有64个FP32 CUDA核心总共6912个核心。RTX 3090则是82个SM每个SM 128个FP32核心总共10496个核心。注意这个“核心”和CPU那个核心不是一个概念CUDA核心更接近CPU里的ALU它的职责就是执行单个浮点运算或整数运算本身没有完整的指令控制能力。关键点在于SM是GPU执行计算的基本单位。你把一个任务丢给GPU驱动和CUDA运行时会把这个任务拆成若干个小块分配到不同的SM上。每个SM内部再把这一个小块进一步拆成线程束warp来调度执行。线程束是GPU调度的真正粒度——32个线程为一组SIMT单指令多线程模式下这32个线程同时执行同一条指令只是各自处理不同的数据。这个设计初衷就是为了让大量核心在同一时刻保持“整齐划一”的忙碌状态避免出现有些核心等指令、有些核心空闲的情况。这套硬件设计对模型推理有几个直接影响。第一矩阵乘法类算子天然适合GPU矩阵乘法本质是大量独立的乘加运算可以完美地映射到大量并行线程上。第二分支逻辑越少越好同一个线程束里只要有线程走不同的分支这个线程束就会被迫串行执行所有分支路径性能会受到明显影响。所以你在写CUDA kernel时要尽量避免类似if (threadIdx.x % 2 0) {...} else {...}这种按线程ID分叉的逻辑。2.2 从Ampere到Hopper再到Blackwell架构演进对推理的影响选GPU硬件时很多人只看显存大小和价位很少关注架构代际差异。但架构差异对推理性能的影响远比你想象的大。以最近几年最常用的三代架构来说架构代表显卡关键改进对推理优化的直接影响AmpereA100, RTX 30系引入第三代Tensor Core支持TF32和稀疏化FP16矩阵运算吞吐大幅提升TF32可以在不损失太多精度的情况下跑FP32精度的网络HopperH100, RTX 40系第四代Tensor CoreTransformer引擎DPX指令针对Transformer结构做了指令级优化FP8支持让推理显存和带宽压力骤降BlackwellB200, RTX 50系第五代Tensor Core第二代Transformer引擎2倍于Hopper的FP4推理吞吐但兼容性问题也会更突出新一代架构通常会引入新的计算精度格式如TF32、FP16、BF16、FP8、INT8这些格式直接决定了推理时的吞吐上限。比如你用FP16代替FP32做推理A100上的FP16 Tensor Core吞吐是FP32的2倍H100的FP8吞吐又是FP16的2倍。这不是简单的“精度砍半、速度翻倍”而是硬件上Tensor Core本身就是为低精度高并行设计的。所以模型部署里的精度量化本质上是赌硬件的低精度计算单元能带来多大的加速收益。这里有一个很多人忽略的点量化带来的加速收益更多来自内存带宽节省而非计算加速。尤其在Decoder推理自回归生成场景下模型是batch size很小的逐token生成计算强度很低瓶颈几乎全在显存带宽上。把权重从FP32砍到FP8意味着同样时间内能从显存里读出4倍的数据这才是量化在LLM推理中收益巨大的根本原因。2.3 Tensor Core推理优化的关键加速单元Tensor Core是NVIDIA从Volta架构开始引入的专用矩阵乘单元它本质上就是一个“一次算一个4x4或1x4矩阵乘”的专门硬件块。普通CUDA核心算矩阵乘是一个个乘加指令循环执行Tensor Core则是直接把一个小矩阵整体吞进去一次输出一整块矩阵乘结果。为什么这很重要因为深度学习推理里绝大多数计算量都集中在卷积和矩阵乘上。以Transformer模型为例QKV投影、注意力分数计算、输出投影全都是矩阵乘。Tensor Core的出现相当于给这些算子修了一条高速公路。在推理优化时Tensor Core的利用情况是必须盯住的指标。用Nsight Compute看kernel的分析报告时如果你看到Tensor Pipe UtilizationTensor Core利用率只有百分之十几说明kernel没有吃的Tensor Core红利该检查数据排布和分块策略了。通常优化思路有这么几层确保数据精度匹配Tensor Core要求A100支持TF32、FP16、INT8H100/B200额外支持BF16、FP8、FP4尽量用官方库cuBLAS、cuDNN而不是自己写通用矩阵乘自己写kernel时分块大小要对齐Tensor Core的输入尺寸要求通常都是用16x16或16x8的分块我在实际部署中见过一个很典型的案例同一个矩阵乘算子用cuBLAS跑只要0.3毫秒用PyTorch默认的优化版约0.5毫秒而如果有人用基础CUDA kernel没做tiling直接裸写耗时可能直接飙到5毫秒以上。差距有10倍以上原因就是没有利用好Tensor Core和共享内存。3. CUDA编程模型软件怎么编排硬件干活3.1 网格、线程块、线程三级结构到底怎么理解CUDA编程模型的核心是一个三级线程结构网格Grid→ 线程块Block→ 线程Thread。对应到GPU硬件上一个Grid里的所有线程块会被分发到不同的SM上执行一个线程块里的所有线程会在同一个SM内部执行。线程块是软件上最小编排单位SM资源共享内存、寄存器、线程数上限决定了一个SM最多能同时驻留多少个线程块。这里有个经常被误解的概念线程块内部才能用共享内存和block级同步不同线程块之间没法直接同步也不该互相通信。如果你把一个需要跨块协作的算法硬塞进来要么需求本身就是不合理的应该重新设计数据切分要么需要额外用全局内存加原子操作或网格级同步CUDA 9之后有cooperative groups支持但性能代价很大。实际工作中我判断线程组织结构通常就三个标准让每个线程负责的工作量尽量平均让线程块大小是32的倍数32线程1个线程束避免线程束内出现空槽位总线程数不要远大于实际需要的并行工作量避免过度创建线程导致调度开销比如你要在GPU上一个长度100万的向量做逐元素操作最省事的做法是开256个线程的线程块总共3907个线程块。每个线程处理一个元素用gridDim.x * blockDim.x threadIdx.x算出自己的全局索引。块的个数略多于SM的驻留能力没关系GPU的块调度器会自动分配但太多的小块反而增加调度开销。之前调一个向量Add kernel把块大小从128调到256耗时掉了差不多15%原因是块太少时每个SM隐藏访存延迟的能力变弱了。3.2 线程束是性能调优的基本观察单位线程束这个概念我觉得是CUDA性能最核心的东西没有之一。GPU的世界里SM每次都是以线程束为单位取指令、派发指令的。一个线程块里有256个线程SM实际会把它拆成8个线程束每32个线程一个束来调度。硬件上有一个隐藏延迟的机制叫线程束切换。当一个线程束执行内存加载这种高延迟操作时SM会立刻切换到其他准备好的线程束继续执行计算指令等内存数据回来了再切回来。所以对于GPU来说足够的并行度 隐藏延迟的资本。这也解释了为什么GPU跑深度学习模型时batch size越大单样本吞吐越高。batch大→每层算子的矩阵乘规模大→并行度足→SM能通过线程束切换把访存延迟藏得干干净净。反过来batch为1的LLM解码阶段矩阵乘规模其实很小比如一个1x512的向量乘512x512的权重矩阵并行度不足SM的调度器经常空转等待显存返回数据。一种比较有效的优化方案是动态批处理dynamic batching。把同一时刻到达的多个请求攒在一起凑成一个更大的矩阵乘一次喂给GPU。这在LLM推理服务里已经是事实标准了。vLLM的continuous batching就是这么干的。理解了线程束调度原理你就能明白为什么这个设计能带来巨大收益而不是仅仅“看起来更合理”。3.3 写一个最简单的CUDA Kernel需要什么光看概念不够怎么上手写和编译CUDA代码我经过反复实践后认为最快的路线是装好CUDA Toolkit后用nvcc编译一个向量加法程序。__global__ void vecAddKernel(const float* a, const float* b, float* c, int n) { int idx blockIdx.x * blockDim.x threadIdx.x; if (idx n) { c[idx] a[idx] b[idx]; } }编译nvcc -O3 -archsm_80 -o vec_add vec_add.cu这里-archsm_80表示针对Ampere架构A100/RTX 30系优化。如果你是H100或RTX 40系用sm_90B200或RTX 50系用sm_120。不过要注意sm_120目前和很多老库存在兼容性问题如果编译时提示“cuda capable sm_120 is not compatible”要么降级到sm_100要么检查PyTorch/TensorRT的CUDA支持版本。我自己在RTX 5070 Laptop上踩过这个坑最新PyTorch也要到2.6以上才支持Blackwell架构。编译完成后用cudaMalloc分配显存cudaMemcpy把数据从CPU拷贝到GPU执行kernel再拷贝回CPU。这基本是早期学习时最常见的“程序人生显存检查、流管理、内存拷贝”。CUDA C对运行时检查很严格大部分已知的CUDA error都要靠查返回值和处理cudaError_t类型去定位少了这步代码出错时经常会看到随机性的数据错乱很难排查。更贴合实际部署的做法是直接用PyTorch的torch.cudaAPI写自定义kernel借助它的CUDA封装简化掉繁琐的设备内存管理流程。一个PyTorch自定义CUDA算子通常只写一个CUDA C扩展用pybind11和torch.utils.cpp_extension.load_inline机制加载。对刚上手的人来说这个路线比纯C开发简单不少又比纯Python循环快几个数量级。我强烈推荐先用这种方法跑通一个完整流程再回头啃核心里更深入的优化技巧。4. 实操从零优化一个CUDA算子手把手调一遍4.1 环境准备和验证别一上来就踩版本坑CUDA开发和模型部署最大的坑是环境不匹配这句话我必须放在最前面说。CUDA Toolkit、NVIDIA驱动、PyTorch/TensorRT这三者各自对版本有严格约束。日常接到报错私信里超过一半是版本搭配导致的。比较省心的搭配经验是这样先装驱动驱动版本决定你能支持哪些CUDA的最高版本再装CUDA Toolkit可以同时装多个版本习惯上用sudo apt install cuda-toolkit-12-4这种装12.412-8装12.8切换用update-alternatives或设置CUDA_HOME环境变量然后装PyTorch用官方pip install torch --index-url https://download.pytorch.org/whl/cu124这种方式保证PyTorch里的CUDA版本匹配你的Toolkit最后确认一切OKpython -c import torch; print(torch.version.cuda, torch.cuda.is_available())验证GPU驱动的经典命名是nvidia-smi不是nvidia-smi -a。看右上角的CUDA Version它代表当前驱动能兼容的最高CUDA运行时版本并不代表你机器上实际装了那个版本的Toolkit。很多人以为这个数字就是当前CUDA环境版本结果版本对应不上一堆稀奇古怪的报错。如果你要用容器形态部署模型直接在Docker镜像里装CUDA全量Toolkit很浪费镜像体积建议直接用官方的nvidia/cuda:12.4.1-runtime-ubuntu22.04镜像只需要一个运行时就够了里面包含了运行CUDA程序所需的所有动态库libcudart.so、libcublas.so等编译工具链和头文件不包含在内。开发编译阶段才需要devel镜像。4.2 一个具体案例优化向量归约Reduction算子向量归约把一个大数组求和成一个值是很多推理场景的基础算子像LayerNorm、Softmax里的求和、attention分数归一化都离不开它。这个算子看起来简单要调好其实不容易是个特别好的练手对象。基础版本写起来很简单__global__ void sumKernel(const float* input, float* output, int n) { int idx blockIdx.x * blockDim.x threadIdx.x; double sum 0.0; for (int i idx; i n; i gridDim.x * blockDim.x) { sum input[i]; } // 用原子加把当前线程的部分和累加到output if (threadIdx.x 0) { atomicAdd(output, (float)sum); } }这个版本有多慢实测下来长度为1000万的数组求和跑出来大约需要0.8毫秒。听起来不多但放进LayerNorm里网络每一层都要跑好几遍这种操作累积下来对推理延迟影响不小。优化的第一步是树形归约。每个线程块内部先把线程的值规约成一个数再用原子操作做块间合并而不是万万个线程全去抢同一个地址做原子加。原因是原子操作本质上是在做串行化大量线程同时往同一个内存地址累加性能会被锁冲突拖垮。__global__ void sumKernelTree(const float* input, float* output, int n) { __shared__ float shared[256]; int idx blockIdx.x * blockDim.x threadIdx.x; float val (idx n) ? input[idx] : 0.0f; shared[threadIdx.x] val; __syncthreads(); // 树形归约stride思想 for (int stride blockDim.x / 2; stride 0; stride 1) { if (threadIdx.x stride) { shared[threadIdx.x] shared[threadIdx.x stride]; } __syncthreads(); } if (threadIdx.x 0) { atomicAdd(output, shared[0]); } }这里有个很容易被新手忽略的点__syncthreads()必须放在每个线程都执行到的地方否则线程块会死锁。你可以想象成全班一起做操一个人不做完当前动作其他人就得一直等他那个不做动作的同学如果在后面直接跳到了下一轮前面的人就永远等不齐了。树形归约版本实测在同样数据量下可以跑到0.2毫秒左右是朴素版本的4倍。但还可以继续优化——最后一个线程块做最终归约完全省掉原子操作。做法是每个线程块归约完自己的部分和之后把结果写到一个全局数组的概率比如blockDim.x个位置然后让一个单独的线程块接管最终合并。这样原子操作一次都不需要发生性能还能再提15%~20%。这个优化过程能说明一个非常关键的观点写CUDA kernel时性能瓶颈往往是并发控制原子操作、同步而非算术本身。很多人在写kernel时只顾着算数却把所有并发问题都丢给了硬件解决结果性能上不去。4.3 用Nsight工具定位瓶颈而不是靠猜我在做推理优化时最常被问到“你是怎么知道该优化哪里”答案很简单用profiler跑一遍让数据告诉我不是靠猜。NVIDIA官方提供了三件套工具nvidia-smi看实时GPU利用率和显存占用适合做整体观察Nsight Systems全系统级别的性能分析能看到CPU和GPU的时间线、kernel启动间隙、数据传输时间适合定位宏观问题比如CPU预处理太慢导致GPU空转Nsight Compute单kernel级别的微架构分析能看到每个kernel的SM利用率、内存吞吐、线程束状态适合精调kernel内部实现Nsight Compute报告里有一个“Speed of Light”速度之光指标图会显示SM计算单元吞吐和内存吞吐的百分比。如果SM throughput用满了而Memory throughput很低说明计算是瓶颈反过来则说明访存是瓶颈。还有的kernel两边都用不满那大概率是并行度不足或存在同步等待问题。用这个工具看归约kernel你会发现基础版本的主要瓶颈在原子操作串行化而树形归约版本主要瓶颈就转移到了最终一次全局写和线程束分歧上。这时候再用__shfl_down_sync做线程束内归约warp shuffle可以做到一个线程束内通过寄存器直接交换数据完全绕过共享内存进一步提升归约速度。这个优化到位后向量归约的耗时能从0.2毫秒再压到0.1毫秒左右。对你没有看错同一份逻辑从粗放到精细优化能差8倍以上的性能。5. 内存层级与访存优化推理性能的真正胜负手5.1 GPU内存金字塔到底哪一层负责什么很多开发者刚接触GPU时觉得显存就是显存直到第一次发现kernel跑得慢去看profiler才发现问题全出在访存身上。GPU内部的内存层级其实是分得非常明确的每一级的速度和容量差别极大。从快到慢粗略排一下存储层级容量相对速度A100示例作用寄存器文件每个SM 256KB最快几乎无延迟每个线程私有的最高速存储共享内存每个SM 192KB几十TB/s级带宽线程块内共享速度接近寄存器L1/L2缓存L1 256KB/SML2 40MBL2约6~8TB/s自动缓存全局内存数据全局显存HBM数十GB约2TB/s主存储区所有线程可见主机内存数十GB~数百GBPCIe 4.0约32GB/s数据进出GPU的通道这个速度差异比是什么概念全局内存的延迟大概400~600个时钟周期而寄存器访问只需要几个时钟周期差了差不多两个数量级。共享内存虽然快但容量极小一整个SM才几百KB。所以CUDA优化的本质就是在这么狭窄的高速存储空间里想方设法把最常用的数据塞进去把不常用的数据按块搬到全局内存再让所有线程尽量复用。我用一个生活化的例子解释GPU的全局内存像是一个大图书馆书都在馆里但每次去借书都要走一段很长的路几百个周期。共享内存像是研究室的阅览桌你可以把一批常用书搬到桌上几个组员围着这张桌子快速翻阅省去了反复跑图书馆的时间。5.2 共享内存Bank Conflict一个让新手一脸懵的坑共享内存虽然是“高速缓存”但它内部又分成32个Bank有32个存储体每次访问周期内32个线程即一个线程束可以同时访问共享内存。理想情况下如果每个线程访问的地址落在不同Bank里这个访问周期就能在一个周期内完成。但如果有两个线程同时访问同一个Bank的不同地址硬件就必须分两次完成——“Bank Conflict”冲突发生性能直接打对折。再具体一点如果32个线程同时访问同Bank的同地址比如都读共享数组的一个元素硬件是支持“广播”的不产生冲突。但如果16个线程访问Bank 0的地址A另16个线程访问Bank 0的地址B硬件就需要分两个周期。代码看起来没毛病性能却莫名衰减一半这种问题不靠Nsight Compute的Bank Conflict指标很难直观感觉到。处理Bank Conflict的经典技巧是填充padding。比如要把一个float data[32][32]的二维数组以行优先方式存储访问时每列刚好落在同一个Bank上就会导致Bank Conflict。改成float data[32][33]每行多一个float行的起始偏移就会错开Bank分布被打散冲突就解决了。这里的本质是通过加1个元素的偏移让本来整齐排列的存储地址在Bank上也能均匀落在不同的位置上。在模型推理里共享内存用得最多的场景就是矩阵乘法的分块优化。把大矩阵切分成小块比如16x16的小矩阵每个线程块负责一块先把数据从全局内存搬到共享内存然后每块数据被反复使用多次。A100上的经典优化中一个16x16的Tile会让每个浮点数被读取大约16次这极大减少了全局内存带宽的压力。这也就是为什么FP16、INT8精度能让推理变快的一个重要原因——数据量少一半搬数据的压力也跟着少一半共享内存能存下更多批次的数据。5.3 融合和访存模式优化部署里的实际收益来源模型部署优化中最落地的两个方向就是算子融合和访存方式优化。算子融合的出发点很朴素GPU里每个kernel的启动和全局内存读写都有固定开销两个相邻算子如果都写中间结果到全局内存、下一个再读整个全局内存等于白白多走两趟PCIe/HBM。把它们融合成一个kernel中间量直接留在寄存器或共享内存里省掉的访存时间非常可观。举一个具体例子在做Transformer推理优化时QK^T算完之后跟softmaxsoftmax又要做一次row-wise max和sum传统做法会有两个中间矩阵写入全局内存再分别读取。如果把这些全部融合到QK^T的kernel尾部直接用寄存器里的局部矩阵做softmax最后只输出P矩阵一次访存量直接砍掉一半。在大batch场景下这个融合能让单层Transformer的延迟降低20%以上。访存模式优化的另一面是合并内存访问coalesced memory access。GPU访问全局内存时一个线程束的32个线程最好访问一段连续的内存地址这样硬件可以把32个4字节数据组合成一笔较大批量的传输完成。如果线程的访问地址是打散的比如按列读一个按行存储的二维数组那么一个内存请求就会退化成多次小传输带宽利用率惨不忍睹。这个在高性能计算里叫“访存合并性”。在写kernel时最典型的提示是让线程ID和连续索引一一对应。比如数据是一个NxK的矩阵行优先存储要按行做处理时应该让线程ID对应列方向连续位置而不是让线程ID取不同行。很多新手一上来把二维网格的threadIdx.x直接当成行索引结果全局内存访问完全不连续性能可能差出十倍。6. 模型部署中的CUDA实战问题与排查经验6.1 常见问题速查表见了别慌这些年在模型部署和CUDA环境上踩过的坑太多这里整理一个高频问题速查表可以直接当作排错手册。现象大概率原因解决办法CUDA error: no kernel image is available for execution on the device编译时的-arch架构和实际GPU不匹配用-archnative或显式指定对应的sm_XX重编比如A100是sm_80RTX 4090是sm_89PyTorch报CUDA out of memory但nvidia-smi显存还有剩余PyTorch缓存分配器把显存预先占用了nvidia-smi看到的是已分配量不是实际使用量用torch.cuda.empty_cache()或设置PYTORCH_CUDA_ALLOC_CONFmax_split_size_mb:128DataLoader多进程导致CUDA显存翻倍每个workser进程都会初始化CUDA context显存被重复占用减少num_workers或者把数据预处理改成主进程内一次完成程序运行一段时间后突然报CUDA Illegal memory access数组越界、野指针或异步执行时的资源竞争在cudaMemcpy和kernel调用后加cudaGetLastError()检查用Compute Sanitizer复盘定位Windows下PyTorch与CUDA版本不匹配安装时用了默认cpu版本或者PyTorch版本太旧torch.__version__确认配cu118/cu124等带CUDA后缀版本同一个模型TensorRT比PyTorch快不了多少算子融合没做透或模型里有大量动态shape导致规划退化检查Builder的优化策略设置尽量固定batch size开启FP16/INT86.2 我踩过最深的三个坑顺便聊聊教训第一个坑是异步执行带来的数据竞争。CUDA里kernel执行和CPU端代码是异步的CPU提交kernel后立刻返回不会等GPU执行完。这看起来很好但如果你在kernel执行过程中去读它上面写的输出缓冲读到的就是脏数据。特别坑的是这种错误不一定是稳定复现的有时候跑10次出现1次排查起来如大海捞针。养成好习惯需要跨设备同步时用cudaDeviceSynchronize()或cudaMemcpy这个本身就是同步的不要依赖于“上一次跑成功了所以这次也会成功”。第二个坑是动态batch下显存碎片化。LLM推理服务经常要动态批处理每次请求大小不一致显存分配模式很乱。运行一天后明明总显存占用不高但申请一批连续大内存时就会失败。一开始我用的是PyTorch默认的分配器后来换了vLLM或TensorRT-LLM这类自带显存管理方案的框架或者自己实现一个简单的显存池问题就基本消失了。对自研服务来说预先分配一块大显存、以page方式管理分配比频繁调用cudaMalloc高效得多。第三个坑比较冷门但遇到的人不少Multi-Process ServiceMPS和并发kernel之间的资源挤压。我在调多路推理并发时为了提升GPU利用率开了MPS结果反而出现了kernel互相干扰、个别请求延迟飙升的情况。后来发现MPS的共享上下文对显存资源分配颗粒度影响不小并发量小的时候效果不明显并发量一旦上去反而会劣化。如果你的推理负载本身就是多路并发独立请求开MPS之前一定要先用关闭状态下的性能基线做对比。6.3 CUDA进阶还有哪些东西在推理场景里值得学当你把基础的线程组织、内存管理、算子融合都吃透了会发现在推理优化这条路上还有几块内容是联系在一起的。CUDA Graphs是我认为最被低估的技术之一。普通方式调用kernel时每个kernel的启动开销大约5~10微秒如果一个推理模型有100个算子光启动开销就0.5~1毫秒。在追求低延迟的场景这个开销占比很可观。CUDA Graphs的作用是把整个模型的计算图提前捕获提交一次就把一串kernel按依赖关系全部发射出去单次推理的启动开销直接被压缩到一个很低的水平。TensorRT里有专门的enqueueV3接口支持图模式PyTorch 2.x也支持torch.cuda.graphs拿来做推理延迟优化效果非常直接。**持久化kernelPersistent Kernel**则是另一种思路。常规kernel执行完就释放资源下一算子又要重新分配和调度。在推理流水线里可以让一个kernel一直驻留GPU循环处理多个请求减少重复的初始化开销。这种模式在正式的推理引擎里用得很多缺点是编程复杂度高、调优难度大建议先从CUDA Graphs入手打牢基础后再进阶。最后一定要提的是精度与数值稳定性问题。位码优化调得太猛显存压力和速度是上去了但模型精度掉了一个多点用户立刻感知到回答质量问题这在实际业务里属于重大事故。在推理部署时所有的精度压缩都要经过离线评测集验证并且要定义好“误差阈值”别只看损失函数直接拿线上真实数据分布来测试。写在最后的一点体会从最开始在PyTorch里直接掉kernel到后来自己维护一个推理引擎中间被无数隐晦的性能问题折磨过。现在回头看CUDA和GPU架构知识真正的价值不是让你丢掉PyTorch自己从头写所有算子而是当框架自带的优化手段不够用时你能清楚地知道瓶颈在哪个环节、用什么手段去解。我最常推荐给团队同事的思路是先用profiler把最耗时的三个算子列出来只看不猜然后判断瓶颈是访存还是计算看Speed of Light最后针对性地做算子融合、分块优化或精度调整。这套方法对新手和进阶者都适用本质上是一种“用数据找真相”的工作方式。如果你正准备把模型部署和推理优化当长期方向CUDA这门基本功值得花一年半载慢慢啃不需要背常用API但要把线程组织、内存层级、调度模型这几块吃透。地基不牢的优化基本靠运气地基牢了之后你会在心里慢慢形成一张“性能地图”拿到什么模型都能快速定位瓶颈所在。这是没有捷径、但足够扎实的路子。
