我搞了快十年的 GPU 计算从 CUDA 3.2 时代一路折腾到现在基于 Hopper、Ada Lovelace 架构的各种加速卡最大的感受是太多人把 GPU 当成“一块跑得很快的显卡”却很少有人真正去理解它底层那套和 CPU 完全不同的计算哲学。计算机体系结构里最核心的一个概念就是 SIMD而 GPU 把这条思路发扬光大成了 SIMT正是这个设计决定了为什么 GPU 适合做矩阵乘法、适合跑大模型也决定了你在 PyTorch 里随手一行model.cuda()的背后到底发生了什么。这篇文章我想从体系结构的角度出发把 GPU 与 SIMD 的关系彻底讲透然后顺着这条主线延伸到 CUDA 编程模型、算子执行全流程以及日常绕不开的环境配置、资源调度和故障排查。无论你是正在复习体系结构课的学生还是准备入门 GPU 算子开发、微调大模型的工程师都可以把这篇当作一份带实操视角的参考笔记。1. GPU 体系结构全景从 SIMD 到 SIMT 的设计哲学1.1 CPU 与 GPU 的本质差异为什么 GPU 要用另一套玩法理解 GPU第一步不是看它有多少个核心而是看它和 CPU 的设计目标有什么不同。CPU 的目标是低延迟它花了大量晶体管做分支预测、乱序执行、多级缓存就是为了让单条指令尽快跑完让你打开浏览器、运行 Excel、处理各种乱序依赖的代码时响应够快。GPU 的目标则是高吞吐它不再关心某一条指令多快而是关心单位时间内能完成多少条指令。打个比方CPU 像一个全能博士给你解微分方程、写论文、修水管都能干但一次只能干一件事GPU 像一万个小学生单独拎出来哪个都不咋样但让他们同时做一万道口算题总量瞬间碾压博士。GPU 的架构设计就是围绕“大规模并行”来的晶体管预算大部分给了 ALU算术逻辑单元和寄存器文件控制逻辑被压缩到很小的比例。所以你会看到 NVIDIA 的 SMStreaming Multiprocessor流式多处理器里面塞了几百个甚至上千个计算单元而控制单元只有简单的取指和调度逻辑。这就带来一个关键机制用并行掩盖延迟。GPU 不指望通过缓存把访存延迟降下来而是让大量线程同时在跑一个线程在等显存数据时调度器立刻切到另一个线程执行把等待时间填满。理解了这一点你就能明白为什么 GPU 的存储层次设计和 CPU 差别这么大。1.2 SIMD 的演化从 CPU 指令集到 GPU 的 SIMT 执行模型SIMDSingle Instruction Multiple Data指的是单条指令同时处理多个数据。CPU 上的 MMX、SSE、AVX 都是这套思路一条 AVX-512 指令可以一次处理 16 个单精度浮点数。它的优点是省指令、省带宽缺点是编程复杂你得按向量宽度手动打包数据而且遇到数据量不整除的情况还要处理尾数。GPU 没有直接走 SIMD 的老路而是提出了SIMTSingle Instruction Multiple Threads单指令多线程执行模型。表面上看它和 SIMD 很像都是“一个指令驱动多个数据单元”但本质区别在于SIMD 是数据级并行程序员看到的是一个宽指令SIMT 是线程级并行程序员写的是普通标量代码每个线程处理一个数据元素硬件自动把一组线程组织起来让它们在同一时刻执行同一条指令。为什么 NVIDIA 当年不直接用 SIMD因为 SIMD 的编程体验太痛苦了。让程序员显式管理 128 位寄存器的打包逻辑会让代码变得极其晦涩而且遇到分支发散divergence时向量宽度白白浪费。SIMT 的高明之处在于程序员以线程为单位思考硬件以 warp 为单位执行。你写C[i] A[i] B[i]这种标量代码硬件把 32 个线程绑成一个 warp每条指令同时作用于 32 个线程的数据最终获得的带宽利用率和手写 SIMD 几乎一样编程难度却低了一个数量级。用生活类比再啰嗦一遍SIMD 像是让一个人同时拿五个锅铲炒五道菜你得练出多线程的手臂SIMT 像是雇五个厨师每人一口锅老板拿着大喇叭喊“放油”五个厨师同时放油。“放油”这条指令是单条的但被五个厨师并行执行了——这就是为什么 GPU 叫“单指令多线程”。1.3 GPU 的存储层次与执行模型带宽才是真正的命脉GPU 的存储层次和 CPU 完全不同理解它是入门体系结构的关键寄存器文件每个线程私有访问大约 1 个周期速度最快容量也最贵。共享内存Shared Memory线程块Block内所有线程共享物理上位于 SM 内部延迟约 20~30 周期是程序员手动管理的高速缓存。全局内存Global Memory就是显存所有线程可见延迟高达 400~800 周期容量大但带宽相对有限。常量内存 / 纹理内存有专门的缓存路径适合广播、局部性强的访问模式。这个层次直接决定了 GPU 编程的核心思路CPU 靠缓存自动优化GPU 靠程序员显式地分块、复用、搬运数据。你在 CUDA 程序里看到的__shared__数组、tiling 循环、bank conflict 优化全都是为了一个目标——把访问全局内存的次数压到最低尽量命中寄存器或共享内存。执行模型的骨架则是Grid一个 kernel 启动→ Block/CTA → Warp → Thread一个 kernel 启动时会创建整个 GridGrid 由若干 Block即前面说的 CTA组成每个 Block 被拆分成多个 Warp32 个线程一组Warp 由 SM 上的 Warp Scheduler 调度执行。GPU 上有多个 SM每个 SM 可以同时驻留多个 BlockBlock 被调度到哪个 SM 是硬件动态决定的。这种灵活性让 GPU 天然支持“可扩展”写好的 CUDA 程序不用改换一个 SM 更多的 GPU 就能自动跑满更多并行度。2. 核心概念拆解Warp、CTA 与 Kernel 执行链路2.1 Warp真正的调度单位很多初学者会把“线程”当成 GPU 调度的最小单位这是最大的误区。线程是程序员看到的逻辑概念Warp 才是硬件真正执行和调度的单位。NVIDIA GPU 里一个 Warp 固定包含 32 个线程一次指令发射32 个线程同时执行除非它们走到了不同的分支路径。为什么是 32这有历史原因也是一个工程平衡的结果太小了调度开销大太大了指令发射的粒度过粗、分支发散时浪费更多。AMD 的 wavefront 是 64 线程也说明这个数字没有唯一正确答案。你在写 CUDA 时经常看到blockDim.x 256这个值必须是 32 的整数倍本质上就是为了让 Block 能整除成 Warp避免出现半个 Warp 的尴尬。Warp 带来的一个经典坑是分支发散if (threadIdx.x % 2 0) { // 路径 A } else { // 路径 B }warp 内线程走到不同分支时硬件会串行执行两个分支先执行路径 A此时偶数线程活跃奇数线程被掩蔽再执行路径 B反过来指令数翻倍。我之前写一个粒子碰撞 kernel有个分支导致性能直接掉了 60%改成用掩码运算统一处理后恢复如初。所以 SIMT 编程的黄金法则第一条就是尽量避免 Warp 内分支或者让分支尽可能在 Warp 粒度上对齐。把线程按threadIdx.x / 32分组来处理分支条件而不是按threadIdx.x % 2是实践中非常实用的技巧。2.2 Cooperative Thread Array线程协作的边界CTACooperative Thread Array就是 CUDA 里的 Thread Block很多资料直接叫 Block。为什么特意强调“Cooperative”因为它定义了一个线程协作的合法范围同一个 Block 内的线程可以通过共享内存交换数据可以通过__syncthreads()做同步可以依赖彼此的写入结果。跨 Block 的线程之间不能同步也不能互相假设执行顺序——这是刻意设计的。理解 CTA 的设计动机要和 GPU 的扩展性挂钩。如果 Block 之间可以互相协作、互相等待那硬件就无法自由地把 Block 调度到不同的 SM 上因为一个 Block 可能在等另一个 Block 执行到某个点但它们可能不在同一个 SM 上甚至可能一个在跑、一个在排队。为了让 GPU 具备“随便调度 Block 都能正确”的能力架构师把协作边界限制在 CTA 内部从而换来了计算规模的自由伸缩。所以我见过一个提得特别好的问题“CTA 和 Warp 是什么关系”一句话回答CTA 是程序员视角的协作容器Warp 是硬件视角的执行单位一个 CTA 由若干个 Warp 组成CTA 内部可以共享内存和同步Warp 之间同一 CTA 内通过__syncthreads()同步Warp 的调度则完全由硬件决定。再延伸一句CUDA 9.0 之后引入的 Cooperative Groups 可以在 Grid 层做协作但那是另一套机制了普通 kernel 里别指望它。2.3 Kernel 在 GPU 上的完整执行流程搞清楚了 Warp 和 CTANext 问题是一个 KernelGPU 算子从 CPU 发起到 GPU 上跑完中间到底经过哪些环节我拆给你看数据准备CPU 侧把数据从内存拷贝到显存调用cudaMemcpy(..., cudaMemcpyHostToDevice)。如果用了统一内存管理这一步可以隐式完成但首次访问会有 page fault 迁移开销。Kernel 启动kernelgrid, block, sharedMemSize, stream(args)CPU 端把它封装成一条启动命令经过 CUDA Runtime 和驱动层写入 GPU 的命令队列。分发GPU 的 front-end 解析启动命令把 Grid 按 Block 拆分分发到各个 GPC/TPC/SM。Block 排队每个 SM 的 work distributor 接收分配过来的 Block放入 Block 调度队列。SM 能同时驻留多少 Block取决于寄存器、共享内存、线程数上限三个资源的综合约束。Warp 调度Block 被拆成 Warp 后Warp Scheduler 按优先级从可执行 Warp 列表里选一个发射指令。这里的“可执行”意味着该 Warp 不阻塞在访存、同步、等待依赖上。结果回传Kernel 结束把结果从显存拷回 CPU 内存D2H。一张表总结阶段负责部件关键开销H2D 拷贝PCIe/DMA 引擎带宽受限常见瓶颈Kernel 启动CPU 驱动 GPU front-end约 2~10 微秒小算子尤其昂贵Block 分发GPU 硬件 work distributor通常可忽略Warp 发射SM 内 Warp Scheduler寄存器、共享内存不足会降低驻留数D2H 回传PCIe/DMA 引擎小结果也要走整条链路这里面最容易翻车的两点一是kernel launch 开销在小算子上比计算时间还长所以有了算子融合二是PCIe 带宽远小于显存带宽很多朴素实现的 GPU 程序其实都卡在内存拷贝上。实际调优时我习惯先拿ncuNVIDIA Nsight Compute或者老牌nvprof看一眼 kernel 的执行时间分布确认瓶颈到底在传输还是在计算再决定下一步优化方向。2.4 从架构到编程模型CUDA 与 PyTorch 的映射关系再往上一层看看 AI 框架是怎么踩在这套体系结构上的。你写model.to(cuda)之后PyTorch 实际上做了什么tensor.cuda()底层先是调用 CUDA Runtime 分配显存然后算子执行时经 ATen dispatch 分发到对应的 CUDA kernel。像最基本的torch.matmul落地是 cuBLAS 的 GEMM卷积落地是 cuDNN你自定义的 PyTorch extension 则直接调用自己的 CUDA kernel。这套映射关系解释了为什么大模型训练必须用 GPUC A B这种矩阵乘核心是乘加操作的海量并行。CPU 的 SIMD 宽度再宽一个核一次也只能处理 16 个浮点GPU 一个 SM 上几百个 FP32 单元整个 GPU 几千个内核同时算。更关键的是大模型推理和训练极度依赖显存带宽——参数要从显存里反复读出来算GPU 的 HBM 带宽动辄 TB/s 级别而 CPU 的 DDR5 只有几十 GB/s。顺便说一句现在大模型用的 FlashAttention 这类算子本质也是“摸清 GPU 存储层次”的产物把Q、K、V分块放进共享内存利用片上 SRAM 的高带宽减少对 HBM 的往返访问。它的每一项优化都能在体系结构里找到对应分块是 tiling不落 reserve 是控制共享内存占用重计算是省带宽换算力。搞懂体系结构再看这些算子优化就会豁然开朗。拓展一下“GPU 驱动开发”这个词。驱动层做的事情本质上就是管理这套执行链路分配显存、维护上下文CUcontext、管理 DMA 传输、解析 kernel 二进制SASS/PTX、处理中断和错误恢复比如后面要讲的 Xid 错误。日常做上层应用的人很少直接碰驱动但排查环境问题时绕不开这些概念。3. 实操GPU 环境搭建、监控与计算资源管理3.1 双显卡环境解析Intel UHD NVIDIA RTX 4060 Laptop 的坑热词里那个“显卡有两个 intel uhd graphics 和 nvidia geforce rtx 4060 laptop gpu”是典型笔记本双显卡配置用的是 NVIDIA Optimus 方案Intel 核显负责屏幕显示输出NVIDIA 独显负责高性能计算和渲染。理想情况下系统会根据负载自动切换但现实里经常出各种幺蛾子nvidia-smi里看不到进程但程序确实在用独显计算——因为显示输出走核显独显只做计算不参与显示就没法在nvidia-smi里看到图形进程。某些老程序默认用核显跑性能惨不忍睹需要在 Windows 的“图形设置”里手动指定程序用高性能 GPU。笔记本的“混合图形”模式还会影响 CUDA 编程——cudaSetDevice(0)拿到的可能不是你预期的那张卡。NVIDIA 在 Windows 下强制单 GPU 模式的问题在一些应用比如 ComfyUI里也遇到过它会提示 “forcing single gpu mode”本质是驱动层对多显卡场景的限制。日常使用小建议BIOS 里可以调核显预留显存大小甚至关闭独显但注意核显和独显是两条独立的资源链路关闭核显只影响显示输出不会把显存让给独显。真要压榨性能优先确认驱动版本和系统设置。3.2 GPU 监控从任务管理器到命令行三板斧Windows 10/11 的任务管理器自带 GPU 监控页能看利用率、显存占用、编解码引擎这足够日常应付了。但 Win7 没有这个功能只能靠第三方工具比如 GPU-Z 看传感器、nvidia-smi.exe在命令行下也能输出实时信息老驱动在 Win7 下仍然能用。Linux 和云服务器才是 GPU 计算的主战场我常用的三件东西nvidia-smi一次性快照看显存和利用率、温度、功耗、当前进程。nvidia-smi dmon连续实时监控每个 GPU 的使用率、显存带宽、温度、功耗排查间歇性性能问题时很好用。nvtop交互式界面类似top适合可视化观察进程级别指标。这里必须提醒一个认知误区GPU 利用率 100% 不代表性能好。有时候 kernel 因为访存冲突、分支发散、占用率低利用率很高但效率很低反过来一个高占用率 kernel 可能因为大量短时间操作导致流水线没喂饱。真正的性能分析要看 Nsight Compute 里的 SM 活动率、内存吞吐、指令分发率这些指标。我见过太多人报障说“GPU 打满了”结果打开 ncu 一看访存指令排队半天ALU 闲着——问题完全在访存优化。3.3 PyTorch GPU 环境配置与验证PyTorch 装 GPU 版本算是个高频问题尤其是第一次配环境的人。核心原则保证 NVIDIA 驱动足够新然后装对应 CUDA 版本的 PyTorch wheel 即可。PyTorch 官方 wheel 自带 CUDA Runtime 和 cuDNN你不需要单独装整套 CUDA Toolkit 也能跑起来除非你要编译自定义算子。最快的安装路径# 创建环境 conda create -n torch python3.10 -y conda activate torch # 安装 CUDA 12.1 对应的 PyTorch pip install torch torchvision --index-url https://download.pytorch.org/whl/cu121验证是否拿到 GPUimport torch print(torch.cuda.is_available()) # True print(torch.cuda.get_device_name(0)) # 例如 NVIDIA GeForce RTX 4060 Laptop GPU print(torch.cuda.mem_get_info()) # (free_bytes, total_bytes) # 简单跑一次矩阵乘法验证 a torch.randn(1024, 1024, devicecuda) b torch.randn(1024, 1024, devicecuda) c a b print(c.sum().item())这个验证能覆盖 90% 的“装完还是用 CPU”的排查路径。如果torch.cuda.is_available()返回 False按优先级检查驱动版本太低或安装损坏——更新驱动。装成了 CPU 版 torch——用pip show torch查看版本号里有没有cu后缀。系统里有多个 Python 环境当前环境没装对。WSL 模式下需要装 Windows 驱动而不是 Linux 驱动。新架构卡比如 RTX 5070 的 sm_120需要对应新版本 CUDA/驱动老工具链会直接报 “is not compatible”。顺便提一下 PaddlePaddle 的验证方法paddle.utils.run_check()会打印是否支持 GPU或者先paddle.is_compiled_with_cuda()看编译选项再paddle.device.cuda.device_count()看设备数量最后跑一个模型试试。原理和 PyTorch 验证类似关键是确认装的是 GPU 版而不是 CPU 版。3.4 GPU 计算资源分配与云原生调度现在很多公司不再自购机器而是租 GPU。云厂商的 GPU 实例按小时计价A100/H100 这种顶级卡和 RTX 4090 这种消费级卡价格差了好几个数量级。选什么卡取决于你跑什么大模型微调吃显存和带宽优先 H100/A100 甚至 H200做 CUDA 开发验证、小规模推理RTX 4090 性价比很高深度造假或者跑 ComfyUI 做生成图也基本吃显存。K8s 里跑 GPU 是个大话题。NVIDIA 官方有 device plugin 让节点上的 GPU 可以以容器资源的形式被调度。为了让一张卡被多个任务共享就有了 MIG多实例 GPU和第三方虚拟化方案。热词里那个 HAMI 就是开源 GPU 虚拟化方案之一它能把物理 GPU 按显存和算力切片分配给不同的 Pod类似“给一块物理卡装上多个独立水龙头”。遇到“GPU 配额已不够预冻结”这种报错通常发生在企业云平台或自建集群的配额系统里提交任务前先冻结一定核时/配额不够就排队等待或申请扩容。这里给个显存估算公式微调时经常用到7B 参数模型FP16 权重占 14GBAdam 优化器的 fp32 主权重再加梯度和优化器状态实际需求大约是权重的 4 倍以上也就是 14GB x 4 ≈ 56GB。所以单卡 24GB 的 3090/4090 必须靠 LoRA 这类参数高效微调方法只训练少量可训练参数或者用梯度检查点和混合精度来压显存。这套估算背后的变量本质都是体系结构里的存储层次和数据类型差异。4. 常见 GPU 问题与排查经验实录4.1 Windows 下的错误代码 43 与驱动崩溃全套流程错误代码 43 大概是 Windows GPU 用户最噩梦的提示了设备管理器里显卡直接显示“Windows 已停止此设备因为其报告了问题代码 43”。这个错误是设备向系统上报了“我启动失败”的信号原因可能来自三方面驱动异常、硬件故障温度/供电/显存、PCIe 链路不稳定。我自己的排查顺序用 DDUDisplay Driver Uninstaller在安全模式里彻底卸载显卡驱动然后重装 NVIDIA 官方稳定版驱动不用 beta 版。打开 Windows 事件查看器找系统日志里的 WHEA 错误和 NVIDIA 相关的 Xid 错误找到具体错误码。如果是新装的驱动出问题回滚版本如果最近 Windows 更新覆盖了驱动用 Windows Update 排除或禁用自动更新驱动。用 GPU-Z 看传感器温度是否异常、PCIe 链路速率是否掉到 x1/x8、显存频率是否异常。以上都排除了那就是硬件问题了。笔记本检查散热、供电台式机重新插拔一次显卡并吹一下 PCIe 插槽。有个容易被忽略的坑部分笔记本双显卡机型在电源模式为“省电”时Windows 可能会让独显直接进入低功耗状态某些老驱动在这种状态下偶发 43。先切到“高性能”模式再重装驱动往往能解决。4.2 Xid 79 与 GPU has fallen off the busXid 79: GPU has fallen off the bus是 NVIDIA 驱动上报的一种错误事件字面意思是 GPU 从 PCIe 总线上“消失”了。驱动对总线上设备的轮询失败于是报告这个错误。常见诱因有供电不足电源老化、线材接触不良、PCIe 链路信号不稳定插槽积灰、弯曲、金手指氧化、过热触发保护、驱动 bug 或超频/降压后不稳定。排查套路Linux 下直接dmesg -T | grep -i xid找日志Windows 用事件查看器找 Xid。记录到 GPU-Z 的日志功能跑个十分钟负载看温度和 PCIe 链路是否稳定。笔记本重点检查是否有省电策略把 PCIe 链路降速了去电源选项里关掉“PCI Express 链接状态电源管理”。台式机拿橡皮擦擦一下金手指重新插紧换一条供电线。如果用了 MSI Afterburner 超频/降压全部恢复默认。我之前遇到过一个很典型的案例一台推理服务器每隔几小时就 Xid 79排查到最后是 6pin 供电线的压线端子老化换了条电源线后三个月没复发。硬件问题往往比软件问题更隐蔽建议优先怀疑供电。4.3 加速不生效Chrome、ComfyUI、Paddle、Ollama、FunASR 逐个排查不同应用报“GPU 加速不可用”的原因五花八门但排查思路共用一套先确认驱动和系统层面的 GPU 是否正常再确认应用真正调用的是哪张卡。Chrome 输入chrome://gpu如果显示 “GPU not support acceleration”通常是 Intel 核显驱动的新版本把某些 OpenGL/D3D 特性搞坏了试着更新或回滚核显驱动笔记本上还可以手动在 Chrome 的启动参数里加--use-angled3d11绕过问题。ComfyUI 在 Windows 上因为 NVIDIA 驱动限制会强制单 GPU 模式双显卡笔记本要在系统设置里让 ComfyUI 进程锁定到独显否则它可能去跑核显。Paddle 的 GPU 验证就按前面 3.3 的步骤来核心是装paddlepaddle-gpu而不是paddlepaddle。Ollama 默认支持 NVIDIA GPUollama ps能看到模型是否加载在 GPU。它后续版本也支持 Intel GPU但需要配合 IPEX-LLM 或特殊配置纯 CPU 机器它会走编译好的优化路径。FunASR 走 onnxruntime必须单独装onnxruntime-gpu否则即使 torch 是 CUDA 版推理还是会落到 CPU速度差一个数量级。这些应用层排错最后都会收敛到同一个结论确认驱动版本、确认 CUDA/CUDNN 库、确认进程实际运行在哪块设备上然后问题就水落石出了。4.4 GPU 虚拟化与多卡调度避坑HAMi、MIG 这类 GPU 虚拟化和切分方案本质是为多租户场景服务的。它们的技术原理都是基于体系结构层面的隔离MIG 是硬件级别的隔离不同实例有独立的 SM 和显存通道性能干扰小HAMi 这种软件虚拟化则通过驱动层拦截显存分配和 kernel 提交把一整块 GPU 切成多个逻辑设备。实际调度中避坑点显存和算力要分开限制只切显存不切算力一个跑重计算的任务会把整卡的算术单元吃满邻居任务跟着遭殃。注意显存碎片虚拟化切分后的显存分配粒度会影响大模型加载模型太大装不进碎片化的小块显存。配额预冻结pre-freeze在集群调度器里提交任务时会先冻结一部分配额防止任务跑到一半发现资源不够。配额不足就会出现“GPU 配额已不够预冻结”的排队报错处理方法是申请扩容或降低任务规格。我在实际中吃过一次亏一个 Pod 申请了虚拟 GPU 的 “2G 显存”但代码里torch.cuda.set_device(0)写死了物理卡号结果驱动层把虚拟设备映射搞错了程序报错。后来统一改成通过环境变量CUDA_VISIBLE_DEVICES指定设备问题立即消失。5. GPU 生态扩展算子开发与更多应用场景5.1 亲手写一个 Kernel从概念到可运行概念聊再多不如手写一个 kernel。下面是 CUDA 里最经典的向量加法麻雀虽小五脏俱全__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]; } }主函数里这样启动int threads 256; int blocks (n threads - 1) / threads; vecAddKernelblocks, threads(d_A, d_B, d_C, n);几个关键点解释一下idx blockIdx.x * blockDim.x threadIdx.x是 CUDA 线程的全局编号映射Grid 里每个 Block 有blockDim.x个线程Block 编号是blockIdx.x。if (idx n)是在处理“数据量不是线程数整数倍”的情况。这个分支是全路径一致的当idx n时整个 Warp 要么全部进入分支要么全部跳过所以理论上不会带来分支发散的性能损失除非 n 不对齐到 32 的倍数。threads 256是常见选择因为 256 整除 32Warp 数量是整数而且每个 Block 的寄存器/共享内存开销适中SM 驻留 Block 数不会太低。blocks数量一般取能填满所有 SM 的值简单估算SM 数 × 每 SM 最大 Block 数 × 2 左右跑起来再微调。这个例子虽然简单但它是理解“SIMT 线程层次 边界处理”的最小单位。我练习算子开发时通常会在这个基础上加共享内存分块写一个矩阵乘然后观察从朴素版到 tiling 版性能翻几倍——这个过程比读十篇论文都管用。5.2 GPU 应用场景图谱不只是 AI 的专属玩具老有人说 GPU 就是跑 AI 的其实 AI 只是 GPU 众多大场景之一。我整理了一张简化的应用图谱场景关键需求典型硬件偏好科学计算分子动力学、气候模拟双精度算力、大显存A100/H100、AMD MI300AI 训练与推理混合精度 Tensor Core、显存带宽H100、RTX 4090、L40S图形渲染VR、游戏、CG光追、光栅化、编解码RTX 游戏卡、专业卡视频编解码NVENC/NVDEC 引擎几乎任何带 NVIDIA GPU 的机器嵌入式边缘计算低功耗、低主频集成 Imagination PowerVR GE8300800MHz这类核显那个“集成 imagination powervr ge8300 gpu频率 800 MHz”的热词刚好说明一件事不是所有 GPU 都追求绝对算力边缘设备里的 GPU 更看重能效比和低成本。PowerVR 是一个历史悠久的移动 GPU 架构和 NVIDIA 的 SIMT 路径并不相同但它的设计初衷依然是“更好的并行度换来更低的功耗”。做嵌入式开发时同样要理解底层并行模型只是优化目标从吞吐变成了功耗。5.3 体系结构知识到底学来干什么很多人问我又不做驱动开发也不写算子学 GPU 体系结构有什么用我的答案很直接它是你做技术决策时的底层地图。当你要部署一个模型时体系结构告诉你显存带宽决定推理吞吐于是你会优先做量化而不是疯狂调 batch size当你写的 CUDA kernel 性能不对时体系结构告诉你可能是访存冲突而不是算力不够当你想选择云上哪种 GPU 实例时体系结构告诉你看 HBM 带宽和核数比例而不只是看显存大小。大学里类似的“GPU 架构与编程”课程其实已经把这条路铺好了先讲计算模型SIMD/SIMT再讲存储层次然后到编程模型和算子优化最后落到实际项目。你要是能把手写 kernel、PyTorch 部署、驱动排错这几个环节都亲自跑一遍这套知识就真正变成了你自己的。我自己做 GPU 性能优化这些年的体会是体系结构的知识往往是在你踩坑之后才真正理解的。第一次遇到 Xid 79第一次被 43 错误折磨第一次发现 kernel 启动开销比计算还大这些经历比任何教科书都让人印象深刻。所以我一直建议新入行的朋友不要只停留在“把模型跑起来”这一步多花点时间去看一个算子到底怎么在 GPU 上被调度、被执行的——你看到的每一条局部性优化、每一次分块计算背后都站着同一个理论框架SIMD 到 SIMT线程到 WarpCTA 到 Grid。把这个框架装进脑子里无论以后做算子开发、模型推理还是云上 GPU 资源管理都会比只背一堆英伟达参数的人走得远得多。