第 1 章 · 导论与环境

学习目标

• 明白 GPU 与 CPU 的根本差异（"很多笨工人" vs "几个聪明工人"）
• 装好 CUDA Toolkit / 或者会用 Google Colab 替代
• 第一次跑出自己 GPU 的 deviceQuery 输出，能解读每一行
• 测出本机的内存拷贝带宽（H2D / D2H / D2D），知道为什么以后要用 pinned memory

前置知识

会用 gcc/clang 编译 C++17 程序；能在终端里运行 make；理解什么是堆栈和指针。不需要任何 GPU / 并行编程经验。

核心概念

1.1 为什么是 CUDA？

大模型推理本质上是一连串巨大的矩阵乘 + 元素级运算。 CPU 一次处理几路（4~64 线程），GPU 一次能并行处理几万路。 以 NVIDIA A100 为例：6912 个 CUDA core + 432 个 Tensor Core， 理论 fp16 算力 312 TFLOPS，是顶级桌面 CPU 的 100 倍以上。

graph LR
    subgraph CPU
      C1["Core × 4-64"] --> Cmem["DRAM
~50 GB/s"]
    end
    subgraph GPU
      G1["CUDA core × 数千"] --> Gmem["HBM 显存
~1-3 TB/s"]
      G1 --> TC["Tensor Core
专攻 MMA"]
    end
    style C1 fill:#f3f1e8,stroke:#8b1538
    style G1 fill:#f3f1e8,stroke:#2f5d3a
    style TC fill:#f3f1e8,stroke:#a86420

关键差距不只在"核数"，更在内存带宽：A100 的 HBM2e 带宽 ~1.5 TB/s，是 DDR5 内存的 30 倍。 LLM 推理是访存瓶颈主导的工作负载（小 batch 下尤其明显），这是为什么 GPU 是必需品。

1.2 环境准备（三选一）

方式 A · Google Colab（推荐初学者，免费）

1. 打开 colab.research.google.com
2. 菜单：Runtime → Change runtime type → Hardware accelerator: T4 GPU → Save
3. 新建 cell，运行 !nvidia-smi，看到 T4 16GB 即成功
4. 克隆本仓库（替换为你的 fork）：

!git clone https://github.com/YOUR/ops.git
%cd /content/ops
!make -C code/ch01_intro ARCH=sm_75 run

方式 B · 本地 Linux + NVIDIA GPU

# Ubuntu 22.04 为例
sudo apt update
sudo apt install -y nvidia-driver-535     # 重启后生效
sudo apt install -y cuda-toolkit-12-4     # 包含 nvcc
echo 'export PATH=/usr/local/cuda/bin:$PATH' >> ~/.bashrc
source ~/.bashrc

# 验证
nvcc --version
nvidia-smi

方式 C · Docker（GPU 已可见）

docker run --gpus all -it -v $PWD:/work \
  nvcr.io/nvidia/cuda:12.4.1-devel-ubuntu22.04 bash
# 容器内：cd /work && make -C code/ch01_intro run

关键代码

1.3 第一个程序：deviceQuery

每个学 CUDA 的人写的第一个程序，是问 GPU 自我介绍。源码： code/ch01_intro/device_query.cu。 关键 API 是 cudaGetDeviceProperties()，它把 GPU 的 100+ 项硬件参数填到一个 cudaDeviceProp 结构体里。

核心代码

cudaDeviceProp p{};
CUDA_CHECK(cudaGetDeviceProperties(&p, /*device=*/0));

std::printf("GPU 0: %s\n", p.name);
std::printf("  SM count           : %d\n", p.multiProcessorCount);
std::printf("  Max threads / SM   : %d\n", p.maxThreadsPerMultiProcessor);
std::printf("  Shared mem / block : %zu B\n", p.sharedMemPerBlock);
std::printf("  Global memory      : %.2f GiB\n",
            p.totalGlobalMem / double(1ull << 30));

// 实测带宽（DDR -> ×2）
double bw = 2.0 * p.memoryClockRate * 1e3 * (p.memoryBusWidth / 8.0) / 1e9;
std::printf("  Peak bandwidth     : %.1f GB/s\n", bw);

编译与运行

cd code/ch01_intro
make ARCH=sm_75    # T4
./device_query

典型输出（T4）

# 请实测并记录本机结果
CUDA driver:  12.2
CUDA runtime: 12.4
Devices:      1

============================================================
GPU 0: Tesla T4
  Compute capability : 7.5  (sm_75, Turing)
  SM count           : 40
  Max threads / block: 1024
  Max threads / SM   : 1024  (= 32 warps)
  Warp size          : 32
  Registers / SM     : 65536 (32-bit)
  Shared mem / block : 49152 B
  Shared mem / SM    : 65536 B
  Global memory      : 14.56 GiB
  L2 cache           : 4096 KiB
  Peak bandwidth     : 320.0 GB/s

逐行解读这些数字：

• SM count = 40：T4 上有 40 个独立的"小处理器"。你写的 kernel 会被切成块，分发到这 40 个 SM 上并行执行。
• Max threads / SM = 1024 = 32 warps：每个 SM 同时驻留最多 32 个 warp（每 warp 32 thread）。GPU 通过快速切换 warp 来"隐藏内存延迟"——这是性能的关键。
• Shared mem / block：以程序查询值为准；shared memory 位于片上，第 5-6 章会专门讲怎么用好它。
• Peak bandwidth：以设备属性与官方规格口径为准。后面用 Nsight 看 roofline 时这是横轴的“屋顶”。

运行结果

1.4 第二个程序：内存带宽实测

源码：bandwidth_estimate.cu。 它做三组对照： pageable H2D（普通 malloc 出来的主机内存 → 显存）、 pinned H2D（cudaMallocHost 的页锁定内存 → 显存）、 D2D（显存内拷贝）。

// pinned (page-locked) host buffer：DMA 不需先复制到内核缓冲
char* h_pinned = nullptr;
CUDA_CHECK(cudaMallocHost(&h_pinned, bytes));

t.start();
CUDA_CHECK(cudaMemcpy(d_buf, h_pinned, bytes, cudaMemcpyHostToDevice));
t.stop();

结果记录表（在目标 GPU 上运行后填写）：

自检清单

练习题

1. cores per SM 表：参考 code/ch01_intro/exercises/01_cores_per_sm_starter.cu，补全函数 cores_per_sm()，让程序能打印 "总 CUDA core 数 = SM 数 × cores/SM"。
   答案见 _solution.cu。
2. 带宽测量：用不同 payload (--MB=1, 4, 16, 64, 256) 跑 bandwidth_estimate，画出 "payload vs 带宽" 曲线。你会发现小 payload 下带宽急剧下降——原因？（提示：固定延迟摊销到的字节数变少。）
3. 查阅：在 CUDA 维基 找出你 GPU 对应的 SM 架构发布年份和 Tensor Core 代际，记在自己的笔记里。

1.7 工业实战：选 GPU、上多卡、看监控

学到这里你能用一张 GPU 跑 hello world。生产环境还要回答四个问题：选哪张卡、怎么上多卡、跑起来怎么监控、用云还是自建。

1.7.1 GPU 选型矩阵（规格核验：2026-07-20）

NVIDIA 有数据中心、专业可视化与 GeForce 等产品线；ECC、NVLink、MIG、散热形态和部署许可应按具体 SKU 与官方条款核验，不能只从架构代号推断。

决策口诀

• 显存优先级 ≥ 算力：LLM 推理大多 memory-bound，显存装不下模型一切归零。70B fp16 ≈ 140 GB，最少 2×A100-80G 或 1×H200。
• 带宽优先级 ≥ 算力：batch 较小的 decode 常受权重与 KV 读取限制，不能只用峰值算力预测 tokens/s。
• 多卡先核验互联：PCIe、NVLink/NVSwitch 拓扑会直接影响张量并行通信，应以目标机器的 P2P 与 collective 实测决定。
• 边缘推理用 L4/T4，功耗低、形态紧凑、TC 算力足够。

1.7.2 多 GPU 互联：NVLink vs PCIe vs IB

训练大模型时多 GPU 间数据交换量巨大（每步几 GB），互联带宽直接决定吞吐。

判断你的卡有没有 NVLink：

nvidia-smi topo -m       # 看 GPU 之间的连接拓扑
# 输出 NV4 / NV8 = NVLink (好); SYS / PHB = 走 PCIe (慢)

多卡通信库永远走 NCCL（all-reduce / all-gather / reduce-scatter / broadcast 等集合通信）。torch.distributed 和 vLLM 的 tensor/pipeline parallel 都基于它。第 8 章的 stream 与第 14 章的扩展任务会用到。

1.7.3 生产监控：nvidia-smi + DCGM

跑生产任务必须监控 GPU 状态。常用命令：

# 实时刷新基本状态 (温度、功耗、显存、利用率)
nvidia-smi -l 1

# 进程级显存占用 + GPU 利用率
nvidia-smi --query-compute-apps=pid,process_name,used_memory --format=csv

# 持续监控的工程化方案 (NVIDIA DCGM)
dcgmi dmon -e 1001,1002,1003,203,150
#  1001 = GR engine active   1002 = SM active   1003 = SM occupancy
#   203 = mem util           150  = mem copy util

# 拉到 Prometheus / Grafana
docker run -d --gpus all -p 9400:9400 nvcr.io/nvidia/k8s/dcgm-exporter
curl localhost:9400/metrics  # OpenMetrics 格式

看 nvidia-smi 时关注这 4 个指标

• GPU-Util：SM 在跑 kernel 的时间百分比。陷阱：100% 不代表算力跑满，它只表示"有 kernel 在跑"。一个 GEMV kernel 也能让 GPU-Util = 100% 但实际 5 TFLOPS。要看 Tensor Active% (Nsight)。
• Memory-Usage：显存占用。高占用可能来自权重、KV cache、workspace 或 allocator reserve；结合请求量与时间序列判断是否泄漏，不能用固定百分比判定。
• Power：vs 卡上限 (TDP)。长期低于 60% TDP 说明在等数据（mem-bound）或在等 launch（小 kernel）。
• Temp：接近具体 SKU 的 thermal limit 时可能降频；同时记录温度、时钟、功耗限制与 kernel 时间，不能套用固定损失比例。

1.7.4 云 vs 自建：建立自己的成本模型

云价随区域、承诺周期、整机 SKU、网络和供给变化，本教程不保留无日期的“每卡每小时”数字。比较时记录供应商 SKU、地区、日期、最短租期、CPU/RAM/网络配额，并把自建的采购、折旧、电力、制冷、运维和闲置率放进同一张表。

1.7.5 环境陷阱 5 连发

1. 驱动版本 vs CUDA Toolkit 版本不匹配：每个 CUDA Toolkit 要求最低驱动版本。nvcc --version 是 Toolkit，nvidia-smi 顶部 "Driver Version: 535.xx / CUDA Version: 12.2" 是驱动支持的最高 CUDA，不是实际 Toolkit。混淆是新手第一大坑。
2. CUDA_VISIBLE_DEVICES 改变物理 ID 顺序：默认按 PCI bus ID 排序，跟 nvidia-smi 顺序不一定相同。多卡训练时容易把同一张卡当作两张。设 CUDA_DEVICE_ORDER=PCI_BUS_ID 强制对齐。
3. MIG (Multi-Instance GPU)：A100/H100 可以切成 7 个独立小 GPU 跑多租户，每片有独立 SM 和显存。但同一时刻只能选"全卡"或"分片"模式。误开 MIG 后程序看到的算力是 1/7，很多人调试半天没头绪。nvidia-smi -mig 0 关闭。
4. 容器里看到的 GPU 数 ≠ 物理数：docker 用 --gpus '"device=0,1"' 限制可见。K8s 用 nvidia.com/gpu: 2 resource。两者都通过 NVIDIA Container Toolkit 实现 isolation。
5. ECC 开启时显存少 ~6%：A100-80G 实际可用 ~76 GB（ECC 校验位占空间）。生产推理通常关 ECC 换更多显存：nvidia-smi -e 0 然后重启。代价：极小概率单 bit 翻转不被检测——不重要的推理可以接受，训练绝对不行。

1.8 研究前沿（2025-2026）

1.8.1 Blackwell 与 Blackwell Ultra：B200、GB200、B300/GB300

2025 年开始 NVIDIA Blackwell 大规模出货，把"训练 / 推理大模型"的硬件单位从单卡升到整柜。

关键概念变化：NVL72 把 72 张 GPU 通过 NVSwitch 全连成"一台逻辑 GPU"，跨卡带宽达 1.8 TB/s（接近单卡 HBM 带宽）。这让 TP=72 成为可能，1.8T 参数模型可以塞进一柜推理。

1.8.2 FP4 / NVFP4 / MXFP4 — 推理新单位

Blackwell 推理硬件原生支持 4-bit 浮点。NVIDIA 力推三种格式：

• NVFP4 (E2M1)：1 sign + 2 exp + 1 mantissa，配 per-block fp8 scale (block=16)。NVIDIA 优选格式，推理精度损失 < 1%
• MXFP4 (E2M1)：开放标准 (OCP MX)，配 E8M0 scale (block=32)。精度略差但跨厂商通用
• MXFP6 (E3M2 / E2M3)：6-bit 中间档，激活用

低精度可减少数据字节并启用不同 Tensor Core 路径，但格式转换、scale、shape、质量约束和 memory/compute bound 会共同决定端到端收益；跨 GPU 的峰值不能直接预测 LLM 吞吐。

1.8.3 模型规模与新范式（2024-2026）

三个大趋势对硬件需求的影响：

1. MoE 主导（DeepSeek-V3、Llama 4、Qwen 3）：参数量 10×，激活 1/3-1/10 → 需要显存大 + Expert Parallelism。NVL72 完美匹配
2. 长 context (1M-10M)：KV cache 显存爆炸 → 必须 PagedAttention + KV 量化 + 长 context attention 优化 (Ring/Stripe)
3. Reasoning models（o1 / R1 风格）：单个回答可能生成 10-100K 个 thinking token → 推理服务的 decode 吞吐变得比 prefill 还关键

1.8.4 推理硬件选型方法（2026）

不要跨厂商比较没有统一 dtype、dense/sparse、功耗和 workload 的峰值。先用模型权重、KV cache、batch/context 算容量，再用 prefill/decode 的 roofline 判断计算或带宽瓶颈，最后在相同质量门槛与 SLO 下测 goodput。采购价格与云价必须绑定供应商、SKU、地区和日期。

1.8.5 学习建议（2026 视角）

• 仍然先学 CUDA，因为 ROCm/MLIR/TPU 接口都是它的"方言"
• 用 H100 或 H200 入门生产级实战；条件允许直接上 B200
• FlashAttention v3 / FlashMLA / ThunderKittens 是 2025 后的必读源码
• DeepSeek-V3 技术报告对理解"原生 FP8 训练 + MoE + MLA + 长 context"四合一极有价值
• 关注 Triton / CUTLASS 3.x CuTe DSL — 未来 5 年的 kernel 开发语言

常见坑

• nvcc: command not found → CUDA Toolkit 没装或没加 PATH。export PATH=/usr/local/cuda/bin:$PATH。
• CUDA driver version is insufficient for CUDA runtime version → 驱动太旧，升级 NVIDIA 驱动到 R535+。
• Colab 跑出来发现没 GPU → 大概率是没切到 GPU runtime，重新 Runtime → Change runtime type。
• 编译报错 unsupported gpu architecture 'compute_xx' → ARCH 参数写错了，按你的 GPU 改 (T4 → sm_75, A100 → sm_80, 4090 → sm_89)。

1.10 CUDA 官方手册精讲（CUDA Programming Guide 13.2（核验：2026-07-20））

从 deviceQuery 到 fatbin 的完整版图

Compute Capability：每张卡的"身份证"

NVIDIA 给每一代 GPU 分配一个 Compute Capability (CC) 号码，形如 X.Y：X 是大版本（架构代），Y 是小版本（同代内的小改）。它直接对应 SM 版本号——CC 8.0 的 GPU 上 SM 就标记为 sm_80。这个数字决定了：

• 有哪些硬件指令可用（FP8 → 需 CC ≥ 8.9；TMA / wgmma → 需 CC ≥ 9.0；FP4 → 需 CC ≥ 10.0）
• 各种资源上限（每 SM 最大 thread、shared memory 大小、register file 容量）
• nvcc 用什么后端编译你的 kernel（-arch=sm_80 意味着"目标 CC 8.0"）

关键规律——"小版本前向兼容、大版本不兼容"：

• cubin 二进制兼容：同一大版本内，更高小版本的 GPU 能加载更低小版本的 cubin（sm_86 可跑 sm_80 cubin），反之不行；大版本之间永远不兼容（sm_86 cubin 在 sm_90 上跑不起来）。
• PTX 前向兼容：PTX 是中间语言，低 CC 的 PTX 可被 JIT 编译到任意更高的 CC。这是同一份二进制能在未来 GPU 上运行的关键。

CUDA 平台栈：你装的到底是什么？

装完 CUDA 之后系统里多了一堆东西。把它们的角色厘清，调试"驱动太旧"或"sm 版本不对"才知道从哪改。

    graph TD
        App["你的 .cu 源码"]
        NVCC["nvcc 编译器"]
        PTX["PTX (虚拟 ISA, 文本)
compute_80"]
        Cubin["cubin (真二进制)
sm_80, sm_86, ..."]
        Fatbin["fatbin 容器
= 多个 PTX + 多个 cubin"]
        Exe["可执行文件 / .so"]
        Driver["NVIDIA Driver
(内核态 + 用户态)"]
        GPU["物理 GPU
(运行 cubin)"]

        App --> NVCC
        NVCC --> PTX
        NVCC --> Cubin
        PTX --> Fatbin
        Cubin --> Fatbin
        Fatbin --> Exe
        Exe -- "运行时加载" --> Driver
        Driver -- "选最匹配的 cubin
或 JIT 编译 PTX" --> GPU

        style App fill:#f3f1e8,stroke:#8b1538
        style GPU fill:#f3f1e8,stroke:#2f5d3a
    

五件套各自的角色

nvidia-smi 显示的 "CUDA Version: 12.6" 到底是什么？

fatbin 的好处：一份二进制走天下

编译时给 nvcc 列一串目标架构，它会把每个架构的 cubin + 一份"最高级 PTX"塞进 fatbin：

# 同时支持 T4 (sm_75) / A100 (sm_80) / H100 (sm_90)
    nvcc kernel.cu \
         -gencode arch=compute_75,code=sm_75 \
         -gencode arch=compute_80,code=sm_80 \
         -gencode arch=compute_90,code=sm_90 \
         -gencode arch=compute_90,code=compute_90 \   # 多塞一份 PTX, 给未来卡 JIT
         -o app

    # 偷懒写法 (CUDA 11.5+): 自动列出所有"主流" CC
    nvcc kernel.cu -arch=all-major -o app

    # 只编当前这张卡 (开发期最快, 但二进制不能移植)
    nvcc kernel.cu -arch=native -o app

运行时驱动按顺序找：① 当前 GPU 有完全匹配的 cubin？用它。② 没有但有同大版本更高小版本的 cubin？也行。③ 都没有但有 PTX？JIT 编译（第一次跑会慢一点，结果缓存在 ~/.nv/ComputeCache）。④ 都没有 → 报 no kernel image is available for execution on the device。

JIT 编译：fatbin 里那份 PTX 在干嘛

fatbin 里塞的 PTX 不是给你看的（虽然能 dump 出来），是为了让程序能跑在 nvcc 编译时还不存在的 GPU 上。流程：

1. 程序启动，Driver 加载第一个 kernel
2. Driver 检查 fatbin：当前 GPU 是 sm_120，fatbin 里没有 sm_120 cubin，但有一份 compute_90 PTX
3. Driver 调用内置的 PTX → cubin 编译器，把 PTX 编成 sm_120 cubin（这就是 JIT compilation）
4. JIT 编译完的 cubin 缓存到 ~/.nv/ComputeCache（Linux）或 %APPDATA%\NVIDIA\ComputeCache（Windows）
5. 后续运行直接读 cache，不再 JIT

控制 JIT 行为的环境变量

NVRTC：运行时编译 CUDA C++ → PTX

PyTorch、Triton、TVM 等框架不会预编译每种张量 shape 的 kernel——它们在运行时用 NVRTC（NVIDIA Runtime Compilation Library）把生成的 CUDA C++ 字符串临时编译成 PTX，再交给 Driver JIT 成 cubin。

// NVRTC 用法骨架 (本教程不深入)
    #include <nvrtc.h>
    const char* src = "__global__ void k() { /* ... */ }";
    nvrtcProgram prog;
    nvrtcCreateProgram(&prog, src, "k.cu", 0, nullptr, nullptr);
    nvrtcCompileProgram(prog, 0, nullptr);          // 编出 PTX 字符串
    size_t ptx_size; nvrtcGetPTXSize(prog, &ptx_size);
    std::string ptx(ptx_size, '\0');
    nvrtcGetPTX(prog, ptx.data());                  // → 传给 cuModuleLoadData() 加载执行

第 14 章的 Mini-LLM 不用 NVRTC（全部预编译），但生产框架不可避免。知道它存在就行。

异构系统的两套内存：host vs device

CUDA 的世界里有两个"主语"，全文都会用到，先固定下来：

    graph LR
        subgraph host["host (CPU 侧)"]
          CPU["CPU cores"] --- HDRAM["host memory
DDR4/5, ~50 GB/s"]
        end
        subgraph device["device (GPU 侧)"]
          SM["SM × N"] --- GMEM["global memory
HBM, ~1-3 TB/s"]
          SM --- SMEM["shared memory
SRAM, on-chip"]
        end
        HDRAM -. "PCIe / NVLink
~16-128 GB/s" .- GMEM
        style host fill:#f3f1e8,stroke:#8b1538
        style device fill:#f3f1e8,stroke:#2f5d3a
    

三种数据搬运

设备代码、主机代码的物理隔离

这是新手最难绕过的一关：

"统一虚拟地址空间" (UVA) ≠ 统一内存

容易混淆的两个名词：

• UVA (Unified Virtual Address Space)：所有 CPU 和 GPU 的内存分配在同一个虚拟地址空间里有唯一地址。这意味着拿到一个指针 p，可以用 cudaPointerGetAttributes(&attr, p) 查出它属于 CPU 还是哪张 GPU。所有现代系统都有 UVA。但 UVA 并不意味着 CPU 能直接读 GPU 内存——只是地址不会冲突。
• Unified Memory / Managed Memory：cudaMallocManaged 分配的内存，CUDA Driver 替你按需迁移到访问者一侧。CPU 访问就搬到 host，GPU 访问就搬到 device。代价：迁移有延迟，性能不可预测。

Warp 与 SIMT：32 个 thread 是一辆"小巴"

"GPU 同时跑数千 thread" 这句话是真的，但 thread 不是一个个独立调度的——它们每 32 个一组打包成 warp，整组一起执行同一条指令。理解 warp 是后面所有性能讨论的前提。

SIMT = "Single Instruction, Multiple Threads"

同一个 warp 的 32 个 thread 共享一个 PC（指令指针）。每个 cycle 整个 warp 一起执行同一条指令，但每个 thread 有自己的寄存器、自己的数据：

cycle t:    所有 32 thread 都在执行 `LDG  r1, [rA + r0*4]`  ← 同一条指令
                但 thread 0 加载 a[0]
                    thread 1 加载 a[1]
                    ...
                    thread 31 加载 a[31]

    cycle t+1:  所有 32 thread 都在执行 `FADD r2, r1, r3`
                但每个 thread 操作自己的 r1, r3, r2

这就是 SIMT——instruction 是 single 的，threads 是 multiple 的，每个 thread 看到的数据不同。和 CPU 上的 SIMD（AVX、SVE）有共通也有区别：

Warp Divergence —— 分支的代价

warp 内 32 个 thread 必须执行同一条指令。如果代码里有 if (cond) { A; } else { B; }，部分 thread 走 A、部分走 B 时怎么办？答案：

所以 GPU kernel 的第一性原理：

1. 同一 warp 的 32 个 thread 尽量走同一条路（写代码时按 warp 边界对齐 if 条件）
2. block 大小总是 32 的倍数（否则最后一个 warp 部分 lane 浪费）
3. 访存按 warp 整体规划（这就是后面要讲的"coalesced memory access"）

第 4 章会用动画演示 divergence，第 5 章会讲 coalescing。这里先记住"warp = 32 thread 同步执行"。

下一章导览

环境就绪后，第 2 章我们写真正的第一个 kernel，理解 <<<grid, block>>> 这个看起来很怪的语法到底在干什么。