这一章想回答的是一个实际的问题：从一张干净的 GPU 开始，到写出一个“跑得好”的 Kernel，中间要经历哪些环节？
我们大致会走三步：
- 环境：先把开发环境搭起来。生产环境主流走容器路线（NVIDIA Container Toolkit + CUDA 镜像，见 环境准备目录）；无容器条件的工作站/教学机可走裸金属安装。
- 范式：先从 CUDA SIMT 的基础入手，然后接触 Streams 带来的异步并发，最后过渡到以 Tensor Core 为中心的 Tile-Based 编程范式。
- 工具：在真正写出代码后，用 nvbandwidth、Nsight Compute / Systems 等工具链来回答“快不快、为什么不快”。
走完这三步，就能独立完成从写 Kernel 到定位带宽/算力瓶颈的过程。
“工欲善其事，必先利其器”。在开始写 CUDA 代码之前，需要先把开发环境整顺。生产环境通常走容器路线（NVIDIA Container Toolkit + CUDA 基础镜像）；工作站/教学机等无容器条件的场景走裸金属安装——两条路线见下方条目与环境目录。
这里有几个经常踩坑的关键点：
- 版本对齐：驱动和用户态库（像
libcuda.so）版本必须匹配，否则会出现 ABI 不一致问题。 - 镜像变体选型：
base/runtime/devel各有适用场景，选错了要么镜像臃肿、要么缺工具链。 - 构建优化：多阶段构建能显著瘦身镜像，同时需要明确指定 CUDA 架构符号（如
sm_80/sm_90/sm_100）以匹配目标 GPU。
这也是 vLLM、TGI、Llama.cpp、DeepSpeed 这类框架能在生产集群上“可复现”的前提——一旦环境一致，高层框架的部署才会真正变得可预测。
- NVIDIA GPU 容器环境：原理与构建指南 - 讲清宿主机与容器的职责划分，以及从驱动安装到实战配置的全过程。
- 大模型训练与推理框架的 GPU 镜像构建深度解析 - 对比 vLLM、TGI、Llama.cpp、DeepSpeed 四大框架的 Dockerfile，详解 CUDA 镜像变体选择与构建优化策略。
- 裸金属 CUDA 环境安装指南 (Ubuntu 22.04) - 不依赖 Docker/SLURM 的裸金属路线：Ubuntu 22.04 上安装 NVIDIA 驱动、CUDA Toolkit 与 NCCL，附环境验收清单。
GPU 编程发展到今天，主流上其实有两种思路，对应两种看待“计算负载”的视角：
- SIMT 模型是经典思路，从线程视角出发，以 Thread 和 Warp 为基本单位组织计算。
- Tile-Based 模型更贴近现代硬件，从数据块的视角出发，以 Block / Tile 为单位组织计算，直接面向 Tensor Core 这类现代张量加速单元。
两种范式并非互斥，而是从不同抽象层级描述同一块 GPU 上的计算。
3.1 CUDA 编程基础
不管你最终写的是 PyTorch 算子、vLLM 插件还是自定义核，CUDA 都是 NVIDIA GPU 上几乎所有深度学习框架与高性能算子的底层抽象，也是理解 GPU 行为的第一道关。
入门阶段主要要搞懂两件事：
- 执行层次：Grid / Block / Warp / Thread 这四级结构有什么用、怎么落到硬件上；
- 异步并发：如何用 Streams 让计算与 H2D/D2H 数据传输重叠，避免 GPU “空等数据”。
把这两点想通，后面再看任何 CUDA 代码基本都不会有理解障碍。
- 核心概念：CUDA 核心原理 | 流处理机制
- 概念解析：GPU 编程导论 | SIMT 到 Tile-Based 编程范式的演进
- 进阶专题：NUMA API、设备属性查询、Stream 并发、Launch 开销、CUDA Graphs、Reduction、Tensor Core GEMM、GPU 内存管理、GPU 访问系统内存、GPU 显存互访、Shared Memory 与 Bank Conflict、Warp-level Programming、Multi-GPU 编程、异步拷贝与 Pipeline、性能调优方法论、CUDA 调试实战、手搓 Attention 入门 等共 21 篇，详见 CUDA 目录。
Tile-Based 的思路是把数据块（Tile）做为编程单元，把原本需要手写的底层能力封装成更高层的原语——比如 Tensor Core 的 MMA 指令、共享内存的 Swizzle 排布、以及异步拷贝（cp.async）。
实际的收益很直接：手写高性能 GEMM 或 FlashAttention 这类算子时，之前要难做对的一大堆底层细节，现在能被编译器帮你处理掉一大半，工程门槛明显降低。
- TileLang 快速入门
- 57 行代码的 FlashAttention：TileLang 前向算子源码级拆解 - online softmax 五状态递推、mask 预填与累加语义的配合，附 H100 实测（248/360 TFLOPs）
- 把权重压到 4-bit：TileLang 反量化 GEMM 源码级拆解 - 寄存器与 shared memory 两条反量化路径、LOP3/twiddling 快速解码、wgmma 发射形态实测验证
4. 性能分析与优化
GPU 性能分析不宜只看一个指标，实践上一般从三个视角交叉印证：
- Nsight Compute 看单个 Kernel：SM occupancy、L1/L2 命中率、memory throughput 等指标呈现中微观的执行质量。
- Nsight Systems 看整个系统：CPU 和 GPU 的时间线、NCCL 集合通信的调度和等待，看谁在等谁。
- nvbandwidth 实测带宽：验证 HBM 和 PCIe 实测带宽是否接近理论上限，以排除硬件或拓扑层面的瓶颈。
带宽上不去时，瓶颈常在硬件侧：PCIe 拓扑与 ACS、NVLink 状态、GPUDirect 路径的诊断方法见 硬件架构模块（PCIe / NVLink / GPUDirect / DMA 各一节）。
- nvbandwidth：用法、原理与测试项详解 - 测量 GPU 显存与 PCIe 传输带宽，并解读各项测试的适用场景。
- GPU 间数据传输方法实测 - 4 种 GPU→GPU 传输方法 A100 实测，从 P2P (249 GB/s) 到 UM (2.1 GB/s)，差距 119×。含交互式可视化。
- Profiling 示例代码 - PCIe/HBM/传输效率的可运行 benchmark 集合。
- 其余 7 篇（PCIe 带宽实测、HBM 带宽、功耗管理、传输效率曲线、ncu/nsys CLI、P2P 带宽）见 性能分析目录。
这里的资料大致按 快速入门 → 进阶实战 → 权威参考 三个梯度组织，适合不同阶段按需取用：
- 官方权威：NVIDIA CUDA C++ Programming Guide，词典式查阅的第一去处。
- 中文教材：樊哲勇《基础与实践》、Professional CUDA C，适合系统性读。
- 优化范例：CUDA-Learn-Notes 收集了 200+ 个 Tensor Core / CUDA Core 的极致优化内核，适合“照着拆解”。
- 活跃社区：GPU Mode、CUDA Reading Group 等，能看到最新的工程讨论。
进阶阶段的重心通常放在 Tensor Core MMA、CuTe 布局、寄存器复用，以及 Warp-Specialized 流水这类深度优化技巧上，适合已经熟悉 CUDA 基础、想再抓一点性能的开发者。
- CUDA-Learn-Notes - 涵盖 200+ 个 Tensor Core/CUDA Core 极致优化内核示例 (HGEMM, FA2 via MMA and CuTe)。
以 NVIDIA 官方 Programming Guide 为权威基准，再配上经典中文教材、开源范例仓库和学术/社区讲座，如下几类资源可以按需组合，形成自己的参考矩阵。
书籍与文档：
- 《CUDA C++ Programming Guide》 - NVIDIA 官方权威指南
- 《CUDA C 编程权威指南》 - 经典教材 (Professional CUDA C Programming)
- 《CUDA 编程：基础与实践 by 樊哲勇》 - 中文经典实战教程
- 《CUDA 编程简介: 基础与实践 by 李瑜》
- 《CUDA 编程入门》 - 改编自北京大学超算队 CUDA 教程讲义
代码仓库与示例：
- Nvidia 官方 CUDA 示例 - 官方标准范例库
- 书中示例代码 (Professional CUDA C)
- 学习笔记 (CudaSteps)
- 示例代码 (CUDA_Programming)
- Multi GPU Programming Models - 多卡编程模型示例
社区与讲座：