1. 为什么GPU概念是AI开发者的必修课1.1 从一次推理延迟排查说起去年帮朋友排查一个本地部署的推理服务现象很典型模型加载正常单条请求响应也还行但并发一上来延迟就爆炸GPU利用率却始终在30%上下晃悠。他第一反应是“显卡不够强”准备再买一张。我让他先把nvidia-smi的采样间隔调到100毫秒盯着看了一分钟发现显存占用几乎顶满而SM流多处理器的活跃度很低。问题根本不在算力而在显存带宽和KV Cache的分配策略上——他把max_model_len设成了模型支持的最大值导致每个请求都预留了一大块显存实际并发数被显存卡死计算单元反而在等数据。这件事让我意识到一个普遍现象很多AI开发者能调通transformers、能跑通vllm的demo但一旦遇到性能问题、显存问题、多卡问题就完全不知道从哪个层面下手。原因很简单——GPU的底层概念没有吃透。你知道batch_size调大能提升吞吐但不知道为什么调到某个值之后反而变慢你知道fp16比fp32省显存但不知道什么时候该用bf16、什么时候该用int8你知道多卡要开tensor_parallel但不知道通信开销什么时候会吃掉全部收益。这篇内容就是冲着这个问题来的。我会把GPU从硬件到软件、从单卡到多卡、从训练到推理的55个核心概念拆开讲清楚重点不是背定义而是让你在遇到问题时知道该看哪个指标、该调哪个参数、该避开哪个坑。适合已经能跑通模型、但想真正搞懂底层逻辑的AI开发者也适合正在做推理服务优化、模型部署的工程师。1.2 这55个概念怎么分类才记得住直接罗列55个名词没有任何意义背完就忘。我习惯把它们分成五层从下往上依次是硬件层、驱动与运行时层、计算与内存层、并行与通信层、推理与部署层。每一层解决不同的问题层与层之间有明确的依赖关系。硬件层是物理基础包括SM、CUDA Core、Tensor Core、显存类型这些。驱动与运行时层是操作系统和GPU之间的桥梁CUDA Driver、CUDA Runtime、cuDNN、NCCL都在这一层。计算与内存层是开发者最常打交道的部分kernel、grid、block、thread、shared memory、寄存器、occupancy这些概念都在这里。并行与通信层处理多卡和多机的问题涉及NVLink、PCIe、AllReduce、张量并行、流水线并行。推理与部署层则是把模型真正跑起来的部分包括KV Cache、PagedAttention、Continuous Batching、量化、投机解码等。这样分层的好处是当你遇到一个具体问题时能快速定位到是哪一层出了状况。比如推理延迟高先看推理层的batching策略和KV Cache管理如果GPU利用率低往计算层看occupancy和内存访问模式如果多卡扩展效率差往通信层看带宽和拓扑。2. 硬件层从芯片结构到显存体系2.1 SM、CUDA Core与Tensor Core到底怎么配合SMStreaming Multiprocessor是GPU的基本计算单元你可以把它理解成一个“小工厂”。一张RTX 4060 Laptop GPU有24个SM一张H100有132个SM。每个SM内部包含多个CUDA Core、Tensor Core、共享内存、寄存器文件、调度器。CUDA Core负责通用的浮点和整数运算Tensor Core专门做矩阵乘加运算是深度学习加速的核心。关键点在于CUDA Core和Tensor Core不是二选一的关系而是协作关系。一个典型的Transformer层里矩阵乘法走Tensor CoreLayerNorm、激活函数、softmax的逐元素操作走CUDA Core。如果你发现Tensor Core利用率很低往往不是因为矩阵乘法少而是因为逐元素操作太多、数据在内存和计算单元之间来回搬运。这里有个实操经验在A100上跑FP16的矩阵乘法Tensor Core的峰值算力是312 TFLOPS但如果你把同样大小的矩阵用CUDA Core做FP32运算峰值只有19.5 TFLOPS差了16倍。所以模型里哪些部分能映射到Tensor Core直接决定了整体吞吐。这也是为什么FlashAttention这类工作重要——它通过重新组织计算顺序让更多操作能在片上完成减少了对显存带宽的依赖。2.2 显存带宽为什么比显存容量更致命很多人选显卡只看显存大小比如“24GB够不够跑70B模型”。但实际跑起来会发现即使显存够速度也可能慢得离谱。原因在显存带宽。显存带宽决定了数据从显存搬到计算单元的速率。以RTX 4090为例显存带宽是1008 GB/s而H100 SXM是3350 GB/s差了3倍多。大模型推理是典型的memory-bound任务每生成一个token都需要把模型权重从显存读一遍。假设模型是70B参数、FP16精度权重占用140GB每生成一个token至少要读140GB的数据。在4090上理论最低延迟是140/1008≈0.14秒也就是每秒最多7个token在H100上理论最低延迟是140/3350≈0.042秒每秒约24个token。这还没算KV Cache和其他开销。所以当你看到“某显卡跑某模型只有几token/s”时先别急着说显卡不行算一下显存带宽的理论上限往往发现已经接近极限了。这时候优化方向不是换更强的计算卡而是用量化降低权重大小或者用投机解码减少大模型的调用次数。2.3 显存类型与ECC的那些坑消费级显卡用的是GDDR6/GDDR6X数据中心卡用的是HBM2/HBM3。HBM的优势不只是带宽高还有功耗低、体积小。但这里有个容易被忽略的点ECC纠错码。数据中心卡默认开启ECC但开启ECC会占用一部分显存带宽和容量。比如A100 80GB开启ECC后可用显存会降到约76GB。更关键的是某些云厂商的实例默认开ECC你在本地测试没开部署到云上发现显存不够排查半天才发现是ECC吃掉的。另一个坑是显存频率的动态调整。笔记本GPU在电池模式下会降频显存带宽直接砍半。如果你在笔记本上做推理测试一定要插电并把电源模式调到“高性能”否则测出来的数据没有参考价值。3. 驱动与运行时层CUDA生态的隐形陷阱3.1 CUDA Driver和CUDA Runtime的区别这两个概念经常被混用但它们在版本管理和兼容性上的行为完全不同。CUDA Driver是显卡驱动的一部分负责和GPU硬件直接通信CUDA Runtime是随CUDA Toolkit安装的库提供开发者调用的API。关键规则是CUDA Runtime的版本不能高于CUDA Driver支持的版本。比如你装了CUDA 12.4的Runtime但驱动只支持到12.2程序会直接报错。反过来Driver版本高于Runtime版本是没问题的向下兼容。实操中常见的坑是用conda装了pytorch-cuda12.1但系统驱动是更老的版本跑起来报CUDA error: no kernel image is available for execution on the device。这时候要么升级驱动要么降级PyTorch的CUDA版本。我一般建议在容器里开发把CUDA Toolkit和驱动解耦宿主机只负责装足够新的驱动。3.2 cuDNN、NCCL、TensorRT各自管什么cuDNN是深度学习的加速库卷积、池化、归一化、激活这些操作都有高度优化的实现。PyTorch和TensorFlow底层都会调cuDNN。但cuDNN的版本和CUDA版本有严格的对应关系版本不匹配会导致性能下降甚至报错。NCCL是多卡通信库负责AllReduce、Broadcast、AllGather这些集合通信操作。多卡训练时梯度同步走的就是NCCL。NCCL的性能高度依赖GPU之间的拓扑结构NVLink和PCIe的带宽差了一个数量级。TensorRT是推理优化引擎它会把训练好的模型做图层融合、精度校准、kernel自动调优生成一个针对特定GPU优化的推理引擎。实测下来同样的模型用TensorRT比原生PyTorch推理快2到5倍但代价是编译时间长、灵活性差模型结构变了就得重新编译。3.3 驱动开发中那些反直觉的细节热词里提到了“gpu驱动开发”这里补充几个实际会遇到的点。GPU驱动不只是让显卡亮机它还管理着命令队列、内存分配、上下文切换。当你跑多个推理进程时驱动会在它们之间调度GPU时间片。如果某个进程占着显存不放其他进程可能直接OOM。另一个细节是WDDM和TCC模式的选择。Windows上默认是WDDM模式这种模式下GPU会参与图形显示显存会被系统占用一部分而且计算任务的调度优先级较低。TCC模式是纯计算模式没有图形显示功能显存全部可用于计算延迟也更稳定。如果你在Windows上做推理服务且不需要用这张卡接显示器切到TCC模式会有明显提升。Linux上默认就是类似TCC的行为所以很多优化在Linux上测不出来差异换到Windows就暴露了。4. 计算与内存层Kernel、线程与Occupancy4.1 Grid、Block、Thread的三层结构CUDA的线程组织是三层Grid包含多个BlockBlock包含多个Thread。执行一个kernel时你指定Grid大小和Block大小GPU的调度器把Block分配到各个SM上执行。这里的关键是Block内的Thread可以通过shared memory和__syncthreads()通信Block之间不能直接通信。所以设计kernel时要把需要频繁交换数据的线程放在同一个Block里。比如矩阵乘法通常把输出矩阵分块每个Block负责一个分块的计算Block内的线程协作加载数据到shared memory然后各自计算。Block大小不是随便设的。太小会导致SM利用率不足太大则可能超出寄存器或shared memory限制导致kernel无法启动。经验值是128到512之间具体要看kernel的资源占用。你可以用cudaOccupancyMaxPotentialBlockSize这个API让CUDA自动推荐一个合适的值。4.2 Shared Memory、寄存器与Occupancy的三角关系Occupancy指的是SM上实际活跃的warp数占最大支持warp数的比例。高Occupancy意味着SM有更多线程可以切换能更好地隐藏内存延迟。但Occupancy不是越高越好因为它和寄存器、shared memory是竞争关系。每个SM的寄存器和shared memory是固定的。如果你的kernel用了很多寄存器每个Block能容纳的线程就少Occupancy就低。反过来如果你强行限制寄存器用量可能会导致寄存器溢出到local memory反而更慢。实操中我一般先用--ptxas-options-v看kernel的寄存器和shared memory用量然后算一下理论Occupancy。如果低于50%再考虑优化。但有些kernel天生就是低Occupancy高ILP指令级并行的比如某些GEMM实现这时候强行提高Occupancy反而会降低性能。4.3 内存访问模式合并访问与Bank Conflict全局内存访问最怕的是非合并访问。GPU的全局内存是以128字节为单位传输的如果一个warp的32个线程访问的地址分散在多个128字节段里就需要多次传输带宽利用率直线下降。理想情况下warp内线程访问连续地址一次传输就能搞定。Shared memory的坑是Bank Conflict。Shared memory分成32个bank每个bank宽度4字节。如果warp内多个线程访问同一个bank的不同地址就会发生冲突访问被串行化。经典的解决办法是padding比如把[32][32]的数组改成[32][33]让每行的起始地址错开避免冲突。这些细节在写自定义kernel时至关重要但如果你只是用PyTorch和vllm这些优化已经被框架做掉了。不过理解这些概念能帮你判断为什么某个操作在GPU上慢是计算瓶颈还是内存瓶颈。5. 并行与通信层多卡扩展的收益与代价5.1 数据并行、张量并行、流水线并行的适用场景数据并行是最简单的每张卡有一份完整的模型副本各自处理不同的数据batch梯度通过AllReduce同步。优点是实现简单缺点是每张卡都要存完整模型显存利用率低。70B模型用FP16存要140GB单卡放不下数据并行就没法用。张量并行是把单个矩阵乘法拆到多张卡上。比如一个[4096, 4096]的权重矩阵切成4份每张卡存[4096, 1024]。计算时每张卡算一部分然后通过AllReduce汇总。优点是能跑单卡放不下的模型缺点是通信频繁对带宽要求高。NVLink的卡间带宽是900GB/sPCIe 4.0 x16只有64GB/s差了14倍。所以张量并行在NVLink机器上效果好在PCIe机器上可能负优化。流水线并行是把模型的不同层放到不同卡上数据像流水线一样依次经过各卡。优点是通信量小缺点是存在流水线气泡需要micro-batch来填充。实际部署中通常是张量并行和流水线并行混合使用比如8卡机器上用4路张量并行加2路流水线并行。5.2 AllReduce、AllGather、ReduceScatter的通信量计算AllReduce是数据并行的核心操作每张卡把自己的梯度发出去收到所有卡梯度的和。Ring AllReduce的通信量是2 * (N-1) / N * 数据量N是卡数。当N很大时通信量趋近于2 * 数据量。这意味着数据并行的通信开销和卡数关系不大主要取决于模型大小。AllGather是每张卡收集所有卡的数据通信量是(N-1) / N * 数据量。ReduceScatter是AllReduce的一半先做reduce再scatter。ZeRO优化就是把AllReduce拆成ReduceScatter和AllGather让每张卡只存一部分优化器状态降低显存占用。实测中8卡A100用NVLink做AllReduce70B模型的梯度同步时间大约在几十毫秒级别。如果用PCIe同样的操作可能要几百毫秒训练吞吐直接砍半。所以多卡训练前一定要确认卡间互联是NVLink还是PCIe。5.3 通信与计算的Overlap怎么做多卡训练的理想状态是通信和计算完全重叠卡在算当前batch的梯度时上一batch的梯度同步已经在后台进行了。PyTorch的DDPDistributedDataParallel默认会做这种overlap但需要满足几个条件梯度分桶gradient bucketing要合理bucket太小通信频繁太大则overlap窗口不够NCCL的流优先级要设置正确否则通信可能被计算阻塞。一个常见的调优参数是bucket_cap_mb默认是25MB。如果你的模型层很大可以适当调大减少通信次数。但调太大又会导致显存峰值上升因为要等整个bucket的梯度都算完才能开始通信。我一般会从25MB开始逐步调到50MB或100MB看吞吐变化。6. 推理与部署层从KV Cache到Continuous Batching6.1 KV Cache为什么是推理显存的大头自回归生成时每生成一个token都需要用到之前所有token的Key和Value。如果每次都重新计算计算量会随序列长度平方增长。KV Cache的思路是把之前算过的Key和Value存下来生成新token时直接复用。KV Cache的大小计算公式是2 * batch_size * num_layers * num_heads * head_dim * seq_len * dtype_size。以LLaMA 7B为例32层、32头、head_dim128、FP16单条序列长度2048时KV Cache大约是2 * 1 * 32 * 32 * 128 * 2048 * 2字节约1GB。如果batch_size是16就是16GB。这还没算模型权重本身占的显存。所以推理服务的并发数往往不是被计算能力限制而是被KV Cache的显存占用限制。vllm的PagedAttention就是来解决这个问题的它把KV Cache分成固定大小的block按需分配避免了为每个请求预留最大长度的浪费。实测下来同样的显存vllm能支持的并发数比原生HuggingFace实现高2到4倍。6.2 Continuous Batching与静态Batching的区别静态Batching是等一批请求凑齐了一起送进模型等所有请求都生成完了再返回。问题是不同请求的生成长度不一样短的早就结束了但GPU还得等长的那个跑完期间短请求占用的计算资源白白浪费。Continuous Batching也叫iteration-level scheduling是每生成一个token就检查一次有没有请求完成了完成了就把它踢出去腾出的资源给新来的请求。这样GPU始终在处理有效请求吞吐能提升数倍。vllm和TensorRT-LLM都默认用这种方式。这里有个实操细节Continuous Batching下batch_size是动态变化的你不能用固定的batch_size去估算显存。vllm的gpu_memory_utilization参数控制显存预分配比例默认0.9。如果你的服务还有其他进程用GPU要适当调低否则会OOM。6.3 量化、投机解码与推理引擎选型量化是降低显存占用和提升推理速度的直接手段。FP16到INT8模型大小减半显存带宽需求也减半理论速度提升接近2倍。但量化会带来精度损失尤其是激活值量化对异常值敏感。实践中权重用INT8、激活用FP16的W8A16方案比较稳精度损失小加速效果也不错。投机解码Speculative Decoding是用一个小模型先草拟几个token然后用大模型一次性验证。如果小模型猜对了就省去了大模型的多次前向。实测在代码生成、翻译这类确定性强的任务上加速比能到2到3倍。但小模型的选择很关键分布差太远的话接受率低反而增加开销。推理引擎选型上vllm适合快速部署和灵活调整社区活跃新模型支持快TensorRT-LLM适合追求极致性能的生产环境但编译和调试成本高llama.cpp适合CPU和低资源场景量化方案成熟。我一般先用vllm做原型验证确定参数后再考虑是否迁移到TensorRT-LLM。7. 常见问题与排查技巧实录7.1 GPU利用率低但显存满了怎么办这是推理服务最常见的症状。排查顺序是先看nvidia-smi的显存占用和GPU利用率如果显存接近100%而利用率低于50%基本可以确定是KV Cache或模型权重占满了显存导致没有足够的空间做计算。解决办法有三个方向一是降低max_model_len减少每个请求的KV Cache预留二是启用量化把模型权重从FP16降到INT8三是调整gpu_memory_utilization让vllm更激进地利用显存做batching。我一般先调max_model_len因为很多场景下实际序列长度远小于模型支持的最大值。7.2 多卡训练速度不升反降的排查清单多卡比单卡慢通常不是GPU的问题而是通信或配置的问题。按以下顺序排查排查项检查方法常见问题卡间互联nvidia-smi topo -mPCIe代替NVLink带宽不足NCCL版本python -c import torch; print(torch.cuda.nccl.version())版本过旧不支持新拓扑batch_size检查每卡batch_size多卡后每卡batch太小计算效率低数据加载看DataLoader耗时CPU预处理成为瓶颈梯度同步用profiler看AllReduce耗时bucket设置不合理overlap失败我遇到过最隐蔽的一次是数据加载用了num_workers0单卡时数据加载和计算还能勉强重叠多卡后数据加载完全跟不上GPU大量时间在等数据。改成num_workers4后吞吐直接翻倍。7.3 模型加载报CUDA Out of Memory但显存明明够这种情况通常是显存碎片化导致的。PyTorch的缓存分配器会预留显存但预留的块可能不连续导致大块显存分配失败。解决办法是设置PYTORCH_CUDA_ALLOC_CONFmax_split_size_mb:128限制分配器的最大分片大小减少碎片。另一个原因是模型加载时先加载到CPU再转到GPU如果CPU内存不够会触发swap看起来像GPU OOM。检查free -h看CPU内存和swap使用情况。如果是这个问题可以用low_cpu_mem_usageTrue参数让模型逐层加载到GPU。7.4 推理结果不稳定或精度下降的排查量化后精度下降是预期内的但如果下降太多要检查量化校准集是否具有代表性。用一小段通用文本做校准和用领域内文本做校准结果可能差很多。另外INT8量化对异常值敏感可以用AWQ或GPTQ这类考虑激活分布的方法比朴素的min-max量化稳。如果是FP16下精度就不对检查是否有算子不支持FP16而回退到了FP32导致类型不一致。PyTorch的torch.autocast可以自动处理混合精度但某些自定义算子需要手动指定。8. 我个人在实际操作中的体会折腾GPU这些年最大的体会是大部分性能问题不是靠换硬件解决的而是靠理解瓶颈在哪里。我见过太多人一遇到慢就想着升级显卡结果换了卡发现还是慢因为瓶颈在数据加载、在通信、在KV Cache管理根本不在计算单元。另一个体会是概念要成体系地学不要零散地背。你单独理解“Occupancy”没用得把它和寄存器、shared memory、内存延迟放在一起看你单独理解“张量并行”也没用得把它和NVLink带宽、AllReduce通信量、模型结构放在一起看。这55个概念之所以重要不是因为它们各自多难而是因为它们共同构成了一个完整的认知框架让你在遇到问题时能快速定位到正确的层面。最后分享一个实用习惯每次部署新模型或新服务前先用nvidia-smi dmon跑一分钟记录GPU利用率、显存占用、功耗、温度的基线。出问题时对比基线能快速判断是配置问题还是负载问题。这个习惯帮我省了无数次盲目排查的时间。
GPU核心概念55讲：从硬件到推理的AI开发者必修课
1. 为什么GPU概念是AI开发者的必修课1.1 从一次推理延迟排查说起去年帮朋友排查一个本地部署的推理服务现象很典型模型加载正常单条请求响应也还行但并发一上来延迟就爆炸GPU利用率却始终在30%上下晃悠。他第一反应是“显卡不够强”准备再买一张。我让他先把nvidia-smi的采样间隔调到100毫秒盯着看了一分钟发现显存占用几乎顶满而SM流多处理器的活跃度很低。问题根本不在算力而在显存带宽和KV Cache的分配策略上——他把max_model_len设成了模型支持的最大值导致每个请求都预留了一大块显存实际并发数被显存卡死计算单元反而在等数据。这件事让我意识到一个普遍现象很多AI开发者能调通transformers、能跑通vllm的demo但一旦遇到性能问题、显存问题、多卡问题就完全不知道从哪个层面下手。原因很简单——GPU的底层概念没有吃透。你知道batch_size调大能提升吞吐但不知道为什么调到某个值之后反而变慢你知道fp16比fp32省显存但不知道什么时候该用bf16、什么时候该用int8你知道多卡要开tensor_parallel但不知道通信开销什么时候会吃掉全部收益。这篇内容就是冲着这个问题来的。我会把GPU从硬件到软件、从单卡到多卡、从训练到推理的55个核心概念拆开讲清楚重点不是背定义而是让你在遇到问题时知道该看哪个指标、该调哪个参数、该避开哪个坑。适合已经能跑通模型、但想真正搞懂底层逻辑的AI开发者也适合正在做推理服务优化、模型部署的工程师。1.2 这55个概念怎么分类才记得住直接罗列55个名词没有任何意义背完就忘。我习惯把它们分成五层从下往上依次是硬件层、驱动与运行时层、计算与内存层、并行与通信层、推理与部署层。每一层解决不同的问题层与层之间有明确的依赖关系。硬件层是物理基础包括SM、CUDA Core、Tensor Core、显存类型这些。驱动与运行时层是操作系统和GPU之间的桥梁CUDA Driver、CUDA Runtime、cuDNN、NCCL都在这一层。计算与内存层是开发者最常打交道的部分kernel、grid、block、thread、shared memory、寄存器、occupancy这些概念都在这里。并行与通信层处理多卡和多机的问题涉及NVLink、PCIe、AllReduce、张量并行、流水线并行。推理与部署层则是把模型真正跑起来的部分包括KV Cache、PagedAttention、Continuous Batching、量化、投机解码等。这样分层的好处是当你遇到一个具体问题时能快速定位到是哪一层出了状况。比如推理延迟高先看推理层的batching策略和KV Cache管理如果GPU利用率低往计算层看occupancy和内存访问模式如果多卡扩展效率差往通信层看带宽和拓扑。2. 硬件层从芯片结构到显存体系2.1 SM、CUDA Core与Tensor Core到底怎么配合SMStreaming Multiprocessor是GPU的基本计算单元你可以把它理解成一个“小工厂”。一张RTX 4060 Laptop GPU有24个SM一张H100有132个SM。每个SM内部包含多个CUDA Core、Tensor Core、共享内存、寄存器文件、调度器。CUDA Core负责通用的浮点和整数运算Tensor Core专门做矩阵乘加运算是深度学习加速的核心。关键点在于CUDA Core和Tensor Core不是二选一的关系而是协作关系。一个典型的Transformer层里矩阵乘法走Tensor CoreLayerNorm、激活函数、softmax的逐元素操作走CUDA Core。如果你发现Tensor Core利用率很低往往不是因为矩阵乘法少而是因为逐元素操作太多、数据在内存和计算单元之间来回搬运。这里有个实操经验在A100上跑FP16的矩阵乘法Tensor Core的峰值算力是312 TFLOPS但如果你把同样大小的矩阵用CUDA Core做FP32运算峰值只有19.5 TFLOPS差了16倍。所以模型里哪些部分能映射到Tensor Core直接决定了整体吞吐。这也是为什么FlashAttention这类工作重要——它通过重新组织计算顺序让更多操作能在片上完成减少了对显存带宽的依赖。2.2 显存带宽为什么比显存容量更致命很多人选显卡只看显存大小比如“24GB够不够跑70B模型”。但实际跑起来会发现即使显存够速度也可能慢得离谱。原因在显存带宽。显存带宽决定了数据从显存搬到计算单元的速率。以RTX 4090为例显存带宽是1008 GB/s而H100 SXM是3350 GB/s差了3倍多。大模型推理是典型的memory-bound任务每生成一个token都需要把模型权重从显存读一遍。假设模型是70B参数、FP16精度权重占用140GB每生成一个token至少要读140GB的数据。在4090上理论最低延迟是140/1008≈0.14秒也就是每秒最多7个token在H100上理论最低延迟是140/3350≈0.042秒每秒约24个token。这还没算KV Cache和其他开销。所以当你看到“某显卡跑某模型只有几token/s”时先别急着说显卡不行算一下显存带宽的理论上限往往发现已经接近极限了。这时候优化方向不是换更强的计算卡而是用量化降低权重大小或者用投机解码减少大模型的调用次数。2.3 显存类型与ECC的那些坑消费级显卡用的是GDDR6/GDDR6X数据中心卡用的是HBM2/HBM3。HBM的优势不只是带宽高还有功耗低、体积小。但这里有个容易被忽略的点ECC纠错码。数据中心卡默认开启ECC但开启ECC会占用一部分显存带宽和容量。比如A100 80GB开启ECC后可用显存会降到约76GB。更关键的是某些云厂商的实例默认开ECC你在本地测试没开部署到云上发现显存不够排查半天才发现是ECC吃掉的。另一个坑是显存频率的动态调整。笔记本GPU在电池模式下会降频显存带宽直接砍半。如果你在笔记本上做推理测试一定要插电并把电源模式调到“高性能”否则测出来的数据没有参考价值。3. 驱动与运行时层CUDA生态的隐形陷阱3.1 CUDA Driver和CUDA Runtime的区别这两个概念经常被混用但它们在版本管理和兼容性上的行为完全不同。CUDA Driver是显卡驱动的一部分负责和GPU硬件直接通信CUDA Runtime是随CUDA Toolkit安装的库提供开发者调用的API。关键规则是CUDA Runtime的版本不能高于CUDA Driver支持的版本。比如你装了CUDA 12.4的Runtime但驱动只支持到12.2程序会直接报错。反过来Driver版本高于Runtime版本是没问题的向下兼容。实操中常见的坑是用conda装了pytorch-cuda12.1但系统驱动是更老的版本跑起来报CUDA error: no kernel image is available for execution on the device。这时候要么升级驱动要么降级PyTorch的CUDA版本。我一般建议在容器里开发把CUDA Toolkit和驱动解耦宿主机只负责装足够新的驱动。3.2 cuDNN、NCCL、TensorRT各自管什么cuDNN是深度学习的加速库卷积、池化、归一化、激活这些操作都有高度优化的实现。PyTorch和TensorFlow底层都会调cuDNN。但cuDNN的版本和CUDA版本有严格的对应关系版本不匹配会导致性能下降甚至报错。NCCL是多卡通信库负责AllReduce、Broadcast、AllGather这些集合通信操作。多卡训练时梯度同步走的就是NCCL。NCCL的性能高度依赖GPU之间的拓扑结构NVLink和PCIe的带宽差了一个数量级。TensorRT是推理优化引擎它会把训练好的模型做图层融合、精度校准、kernel自动调优生成一个针对特定GPU优化的推理引擎。实测下来同样的模型用TensorRT比原生PyTorch推理快2到5倍但代价是编译时间长、灵活性差模型结构变了就得重新编译。3.3 驱动开发中那些反直觉的细节热词里提到了“gpu驱动开发”这里补充几个实际会遇到的点。GPU驱动不只是让显卡亮机它还管理着命令队列、内存分配、上下文切换。当你跑多个推理进程时驱动会在它们之间调度GPU时间片。如果某个进程占着显存不放其他进程可能直接OOM。另一个细节是WDDM和TCC模式的选择。Windows上默认是WDDM模式这种模式下GPU会参与图形显示显存会被系统占用一部分而且计算任务的调度优先级较低。TCC模式是纯计算模式没有图形显示功能显存全部可用于计算延迟也更稳定。如果你在Windows上做推理服务且不需要用这张卡接显示器切到TCC模式会有明显提升。Linux上默认就是类似TCC的行为所以很多优化在Linux上测不出来差异换到Windows就暴露了。4. 计算与内存层Kernel、线程与Occupancy4.1 Grid、Block、Thread的三层结构CUDA的线程组织是三层Grid包含多个BlockBlock包含多个Thread。执行一个kernel时你指定Grid大小和Block大小GPU的调度器把Block分配到各个SM上执行。这里的关键是Block内的Thread可以通过shared memory和__syncthreads()通信Block之间不能直接通信。所以设计kernel时要把需要频繁交换数据的线程放在同一个Block里。比如矩阵乘法通常把输出矩阵分块每个Block负责一个分块的计算Block内的线程协作加载数据到shared memory然后各自计算。Block大小不是随便设的。太小会导致SM利用率不足太大则可能超出寄存器或shared memory限制导致kernel无法启动。经验值是128到512之间具体要看kernel的资源占用。你可以用cudaOccupancyMaxPotentialBlockSize这个API让CUDA自动推荐一个合适的值。4.2 Shared Memory、寄存器与Occupancy的三角关系Occupancy指的是SM上实际活跃的warp数占最大支持warp数的比例。高Occupancy意味着SM有更多线程可以切换能更好地隐藏内存延迟。但Occupancy不是越高越好因为它和寄存器、shared memory是竞争关系。每个SM的寄存器和shared memory是固定的。如果你的kernel用了很多寄存器每个Block能容纳的线程就少Occupancy就低。反过来如果你强行限制寄存器用量可能会导致寄存器溢出到local memory反而更慢。实操中我一般先用--ptxas-options-v看kernel的寄存器和shared memory用量然后算一下理论Occupancy。如果低于50%再考虑优化。但有些kernel天生就是低Occupancy高ILP指令级并行的比如某些GEMM实现这时候强行提高Occupancy反而会降低性能。4.3 内存访问模式合并访问与Bank Conflict全局内存访问最怕的是非合并访问。GPU的全局内存是以128字节为单位传输的如果一个warp的32个线程访问的地址分散在多个128字节段里就需要多次传输带宽利用率直线下降。理想情况下warp内线程访问连续地址一次传输就能搞定。Shared memory的坑是Bank Conflict。Shared memory分成32个bank每个bank宽度4字节。如果warp内多个线程访问同一个bank的不同地址就会发生冲突访问被串行化。经典的解决办法是padding比如把[32][32]的数组改成[32][33]让每行的起始地址错开避免冲突。这些细节在写自定义kernel时至关重要但如果你只是用PyTorch和vllm这些优化已经被框架做掉了。不过理解这些概念能帮你判断为什么某个操作在GPU上慢是计算瓶颈还是内存瓶颈。5. 并行与通信层多卡扩展的收益与代价5.1 数据并行、张量并行、流水线并行的适用场景数据并行是最简单的每张卡有一份完整的模型副本各自处理不同的数据batch梯度通过AllReduce同步。优点是实现简单缺点是每张卡都要存完整模型显存利用率低。70B模型用FP16存要140GB单卡放不下数据并行就没法用。张量并行是把单个矩阵乘法拆到多张卡上。比如一个[4096, 4096]的权重矩阵切成4份每张卡存[4096, 1024]。计算时每张卡算一部分然后通过AllReduce汇总。优点是能跑单卡放不下的模型缺点是通信频繁对带宽要求高。NVLink的卡间带宽是900GB/sPCIe 4.0 x16只有64GB/s差了14倍。所以张量并行在NVLink机器上效果好在PCIe机器上可能负优化。流水线并行是把模型的不同层放到不同卡上数据像流水线一样依次经过各卡。优点是通信量小缺点是存在流水线气泡需要micro-batch来填充。实际部署中通常是张量并行和流水线并行混合使用比如8卡机器上用4路张量并行加2路流水线并行。5.2 AllReduce、AllGather、ReduceScatter的通信量计算AllReduce是数据并行的核心操作每张卡把自己的梯度发出去收到所有卡梯度的和。Ring AllReduce的通信量是2 * (N-1) / N * 数据量N是卡数。当N很大时通信量趋近于2 * 数据量。这意味着数据并行的通信开销和卡数关系不大主要取决于模型大小。AllGather是每张卡收集所有卡的数据通信量是(N-1) / N * 数据量。ReduceScatter是AllReduce的一半先做reduce再scatter。ZeRO优化就是把AllReduce拆成ReduceScatter和AllGather让每张卡只存一部分优化器状态降低显存占用。实测中8卡A100用NVLink做AllReduce70B模型的梯度同步时间大约在几十毫秒级别。如果用PCIe同样的操作可能要几百毫秒训练吞吐直接砍半。所以多卡训练前一定要确认卡间互联是NVLink还是PCIe。5.3 通信与计算的Overlap怎么做多卡训练的理想状态是通信和计算完全重叠卡在算当前batch的梯度时上一batch的梯度同步已经在后台进行了。PyTorch的DDPDistributedDataParallel默认会做这种overlap但需要满足几个条件梯度分桶gradient bucketing要合理bucket太小通信频繁太大则overlap窗口不够NCCL的流优先级要设置正确否则通信可能被计算阻塞。一个常见的调优参数是bucket_cap_mb默认是25MB。如果你的模型层很大可以适当调大减少通信次数。但调太大又会导致显存峰值上升因为要等整个bucket的梯度都算完才能开始通信。我一般会从25MB开始逐步调到50MB或100MB看吞吐变化。6. 推理与部署层从KV Cache到Continuous Batching6.1 KV Cache为什么是推理显存的大头自回归生成时每生成一个token都需要用到之前所有token的Key和Value。如果每次都重新计算计算量会随序列长度平方增长。KV Cache的思路是把之前算过的Key和Value存下来生成新token时直接复用。KV Cache的大小计算公式是2 * batch_size * num_layers * num_heads * head_dim * seq_len * dtype_size。以LLaMA 7B为例32层、32头、head_dim128、FP16单条序列长度2048时KV Cache大约是2 * 1 * 32 * 32 * 128 * 2048 * 2字节约1GB。如果batch_size是16就是16GB。这还没算模型权重本身占的显存。所以推理服务的并发数往往不是被计算能力限制而是被KV Cache的显存占用限制。vllm的PagedAttention就是来解决这个问题的它把KV Cache分成固定大小的block按需分配避免了为每个请求预留最大长度的浪费。实测下来同样的显存vllm能支持的并发数比原生HuggingFace实现高2到4倍。6.2 Continuous Batching与静态Batching的区别静态Batching是等一批请求凑齐了一起送进模型等所有请求都生成完了再返回。问题是不同请求的生成长度不一样短的早就结束了但GPU还得等长的那个跑完期间短请求占用的计算资源白白浪费。Continuous Batching也叫iteration-level scheduling是每生成一个token就检查一次有没有请求完成了完成了就把它踢出去腾出的资源给新来的请求。这样GPU始终在处理有效请求吞吐能提升数倍。vllm和TensorRT-LLM都默认用这种方式。这里有个实操细节Continuous Batching下batch_size是动态变化的你不能用固定的batch_size去估算显存。vllm的gpu_memory_utilization参数控制显存预分配比例默认0.9。如果你的服务还有其他进程用GPU要适当调低否则会OOM。6.3 量化、投机解码与推理引擎选型量化是降低显存占用和提升推理速度的直接手段。FP16到INT8模型大小减半显存带宽需求也减半理论速度提升接近2倍。但量化会带来精度损失尤其是激活值量化对异常值敏感。实践中权重用INT8、激活用FP16的W8A16方案比较稳精度损失小加速效果也不错。投机解码Speculative Decoding是用一个小模型先草拟几个token然后用大模型一次性验证。如果小模型猜对了就省去了大模型的多次前向。实测在代码生成、翻译这类确定性强的任务上加速比能到2到3倍。但小模型的选择很关键分布差太远的话接受率低反而增加开销。推理引擎选型上vllm适合快速部署和灵活调整社区活跃新模型支持快TensorRT-LLM适合追求极致性能的生产环境但编译和调试成本高llama.cpp适合CPU和低资源场景量化方案成熟。我一般先用vllm做原型验证确定参数后再考虑是否迁移到TensorRT-LLM。7. 常见问题与排查技巧实录7.1 GPU利用率低但显存满了怎么办这是推理服务最常见的症状。排查顺序是先看nvidia-smi的显存占用和GPU利用率如果显存接近100%而利用率低于50%基本可以确定是KV Cache或模型权重占满了显存导致没有足够的空间做计算。解决办法有三个方向一是降低max_model_len减少每个请求的KV Cache预留二是启用量化把模型权重从FP16降到INT8三是调整gpu_memory_utilization让vllm更激进地利用显存做batching。我一般先调max_model_len因为很多场景下实际序列长度远小于模型支持的最大值。7.2 多卡训练速度不升反降的排查清单多卡比单卡慢通常不是GPU的问题而是通信或配置的问题。按以下顺序排查排查项检查方法常见问题卡间互联nvidia-smi topo -mPCIe代替NVLink带宽不足NCCL版本python -c import torch; print(torch.cuda.nccl.version())版本过旧不支持新拓扑batch_size检查每卡batch_size多卡后每卡batch太小计算效率低数据加载看DataLoader耗时CPU预处理成为瓶颈梯度同步用profiler看AllReduce耗时bucket设置不合理overlap失败我遇到过最隐蔽的一次是数据加载用了num_workers0单卡时数据加载和计算还能勉强重叠多卡后数据加载完全跟不上GPU大量时间在等数据。改成num_workers4后吞吐直接翻倍。7.3 模型加载报CUDA Out of Memory但显存明明够这种情况通常是显存碎片化导致的。PyTorch的缓存分配器会预留显存但预留的块可能不连续导致大块显存分配失败。解决办法是设置PYTORCH_CUDA_ALLOC_CONFmax_split_size_mb:128限制分配器的最大分片大小减少碎片。另一个原因是模型加载时先加载到CPU再转到GPU如果CPU内存不够会触发swap看起来像GPU OOM。检查free -h看CPU内存和swap使用情况。如果是这个问题可以用low_cpu_mem_usageTrue参数让模型逐层加载到GPU。7.4 推理结果不稳定或精度下降的排查量化后精度下降是预期内的但如果下降太多要检查量化校准集是否具有代表性。用一小段通用文本做校准和用领域内文本做校准结果可能差很多。另外INT8量化对异常值敏感可以用AWQ或GPTQ这类考虑激活分布的方法比朴素的min-max量化稳。如果是FP16下精度就不对检查是否有算子不支持FP16而回退到了FP32导致类型不一致。PyTorch的torch.autocast可以自动处理混合精度但某些自定义算子需要手动指定。8. 我个人在实际操作中的体会折腾GPU这些年最大的体会是大部分性能问题不是靠换硬件解决的而是靠理解瓶颈在哪里。我见过太多人一遇到慢就想着升级显卡结果换了卡发现还是慢因为瓶颈在数据加载、在通信、在KV Cache管理根本不在计算单元。另一个体会是概念要成体系地学不要零散地背。你单独理解“Occupancy”没用得把它和寄存器、shared memory、内存延迟放在一起看你单独理解“张量并行”也没用得把它和NVLink带宽、AllReduce通信量、模型结构放在一起看。这55个概念之所以重要不是因为它们各自多难而是因为它们共同构成了一个完整的认知框架让你在遇到问题时能快速定位到正确的层面。最后分享一个实用习惯每次部署新模型或新服务前先用nvidia-smi dmon跑一分钟记录GPU利用率、显存占用、功耗、温度的基线。出问题时对比基线能快速判断是配置问题还是负载问题。这个习惯帮我省了无数次盲目排查的时间。