Repository navigation

CUDA 编程 (CUDA Programming)

这个目录覆盖了从零开始系统学习 CUDA 编程的完整路径：执行模型 → 核心架构 → 异步流 → 编程范式 → 内存管理 → 性能调优 → 调试。21 篇文章外加 12 个可运行 benchmark，大部分配有 A100 实测验证。

1. GPU 编程导论

GPU Architecture and Programming — An Introduction：

介绍了 GPU 的分层执行模型：Grid, Block, Warp, Thread。
解释了 SIMT (Single-Instruction Multiple-Threads) 的基本原理。
包含架构图解与核心概念辨析。

2. CUDA 核心详解

深入解析 Nvidia CUDA 核心（CUDA Cores）的硬件架构。
探讨计算单元的组成与工作方式。

3. CUDA 流处理

详细介绍 CUDA Streams 的概念。
讲解如何利用流实现并发执行（计算与数据传输的重叠）。
异步编程模型的基础。

4. SIMT 到 Tile-Based 编程范式

从 SIMT 到 Tile-Based：GPU 编程范式的演进与实战解析
剖析 NVIDIA cuTile 编程模型。
对比传统 SIMT (Thread 视角) 与 Tile-Based (Block/Tile 视角) 的编程思维。
以矩阵乘法 (GEMM) 为例展示 Tensor Core 的抽象与使用。

5. CUDA NUMA API 编程实践

单 GPU 环境下的 NUMA 亲和性管理。
cudaMallocHost 与 NUMA 节点分配策略。
cudaMemAdvise / cudaMemPrefetchAsync 在 Managed Memory 中的应用。
CPU 亲和性绑定 (taskset / numactl) 的最佳实践。

6. GPU 原子操作与 PCIe 能力查询

cudaDeviceGetAttribute 查询 100+ 种底层硬件能力。
PCIe 原子操作支持确认（Inbound/Outbound Atomic）。
Host Native Atomic 与数据中心 GPU 的能力差异。
RTX 5090 关键属性实测。

7. CUDA Streams 并发实战

单 GPU 上 H2D + Kernel + D2H 重叠执行的完整 demo。
4 个 stream 的实测加速比 2.36x（RTX 5090，2 个 async copy engine）。
Nsight Systems 可视化概念图与 stream 最佳实践。

8. Kernel Launch 开销测量

空 kernel launch 延迟实测：2.6 μs（RTX 5090）。
不同 block 数量对 launch 开销的影响。
CPU vs GPU 决策边界与 CUDA Graph 替代方案。

9. CUDA Graphs 编程

将多次 kernel launch + memcpy 合并为一次 graph launch。
覆盖 Stream Capture 与 Manual API 两种创建方式。
生命周期：录制 → 实例化 → 启动 → 更新。
A100 实测：Instantiation ~33 μs, Repeat Launch ~2.3 μs。

10. Reduction：从朴素实现到 Warp Shuffle

GPU 并行编程最经典的教学案例，展示从 0.048ms 到 0.022ms 的逐级优化（2.18× 加速）。
8 个 kernel 变体：interleaved → sequential → warp shuffle → template → Cooperative Groups。
A100 实测 --shmoo 完整性能表。
Warp shuffle 原理与 shared memory / register 性能对比。

11. Tensor Core GEMM 性能实测

运行官方 cudaTensorCoreGemm / bf16TensorCoreGemm，A100 实测 FP16 52.5 TFLOPS、BF16 90.9 TFLOPS。
理论峰值对标（156 TFLOPS dense）与矩阵大小对利用率的影响。
WMMA API 编程模型 + "算力-带宽-算数密度"三角分析。

12. GPU 内存管理——从推理工程师的日常问题出发

以 Linux 概念为类比，系统讲解 GPU 显存管理的核心机制：虚拟内存与碎片化、DMA 传输（pinned vs pageable）、NVLink/PCIe 拓扑、跨进程共享（CUDA IPC）、GPUDirect Storage、大页与 TLB、内存层级。
配套可交互概念图：gpu-memory-visual.html（6 个图层：物理拓扑 / 内存层级 / DMA 路径 / MMU 页表 / 碎片化 / 跨进程共享）。
包含 CPU ↔ GPU 命令对照附录和性能诊断决策树。
机制层的五种 Linux 映射见第 19 篇：GPU 访问系统内存。

13. Shared Memory 与 Bank Conflict

从 Bank 的硬件拓扑出发，系统讲解 Bank Conflict 的产生机制（stride 访问、列优先写入），以及三种消除方法：Padding（空间换时间）、Swizzle（地址 XOR 变换）、Reorder（算法层面重构）。
包含三个实战场景拆解：GEMM 的 tile 宽度为什么必须是 32 的倍数、FlashAttention 中的 XOR Swizzle 应用、Reduction 的 Sequential Addressing 为何天然无冲突。
跨架构迁移注意事项（V100 4B Bank → A100/H100 的 FP16 双通道 Bank 判定差异）。
附带 Bank Conflict 诊断 SOP 和 Python 验证脚本。

14. Warp-level Programming——从 Shuffle 到 Cooperative Groups

系统讲解 Warp 原语体系：Shuffle（__shfl_sync/__shfl_down_sync/__shfl_up_sync/__shfl_xor_sync）、Vote（__ballot_sync/__all_sync/__any_sync）、Match（__match_any_sync/__match_all_sync）、__activemask。
涵盖 Cooperative Groups tiled_partition 现代抽象与 A100 硬件 __reduce_add_sync 加速指令。
A100 实测验证：Warp Reduce、Prefix Sum Scan、Ballot Filter、Match Dedup。

15. Multi-GPU CUDA 编程——从 P2P 到 NCCL 的多卡协作

系统讲解多 GPU 编程体系：多 device 管理（cudaSetDevice）、Peer Access（cudaDeviceEnablePeerAccess）、P2P Memcpy（NVLink 249 GB/s vs CPU 中转 12 GB/s）、跨 GPU Stream/Event 同步。
涵盖 NCCL 集合通信基础（ncclAllReduce 等）、拓扑感知（NVSwitch domain、nvidia-smi topo -m）。
A100 实测验证：Peer Access、P2P 带宽 20.7× 加速比、跨 GPU 同步、NCCL AllReduce。

16. 异步拷贝与现代 Pipeline——从 cp.async 到 memcpy_async

系统讲解异步数据搬运技术：传统同步 Load 的瓶颈、cp.async PTX 指令（SM80+）、Pipeline 原语（cuda::pipeline）、cuda::memcpy_async C++ API。
涵盖 double buffering / triple buffering 的 Shared Memory 布局与 buffer 管理。
实战：BF16 GEMM 的 cp.async loop、FlashAttention 的 tile pipeline、通用 tile-based kernel 的 pipeline 模板。
A100 验证：1-stage sync vs 2-stage double buffer 对比（同步 ld 实现，加速比有限——证明异步 DMA 是 pipeline 收益的核心）。

17. CUDA 性能调优方法论——从 Roofline Model 到 ncu 诊断

系列收官文章，建立分级诊断框架：Roofline Model（天花板分析）→ ncu 三指标（瓶颈定位）→ 6 级逐级优化 Checklist。
案例：朴素 GEMM 从 Level 0 → Level 6 的完整调优过程（规划加速比 ~80×）。
包含优化边界分析（边际收益递减、"够好了"的判断标准）。

18. CUDA 调试实战——从 cuda-gdb 到 compute-sanitizer

系统讲解 CUDA 调试工具链：cuda-gdb（断点、寄存器/Shared Memory 检查、NaN 定位）、compute-sanitizer（越界/race/未初始化检测、内存泄漏）。
cudaError_t 错误处理模式（同步 vs 异步 error、生产环境宏）。
常见 Bug 排查 SOP：NaN、Hang、间歇性错误、内存泄漏。

19. GPU 访问系统内存——MMU、IOMMU、MMIO、mmap、ioremap 五种机制对照

先把 Linux 的五个机制放回各自的层：MMU / IOMMU 是硬件地址翻译，MMIO 是 CPU 访问设备窗口，mmap / ioremap 是软件映射。
再映射到 GPU：GPU 页表 / GMMU、系统 IOMMU、BAR0 / BAR1、CUDA VMM、驱动 ioremap；区分数据面和控制面。
用五条端到端路径串起来：pageable copy、pinned copy、zero-copy、managed memory、peer / coherent。
含 4 张手绘图，以及 Linux 源码、CUDA 文档和 NVIDIA 开源内核模块的源文件索引。

20. GPU 如何访问另一块 GPU 的显存——同进程 P2P、跨进程 IPC 与 VMM Fabric

区分拷贝、映射和 kernel 直接解引用：cudaMemcpyPeerAsync 不会让远端指针变成本地可解引用，cudaIpcOpenMemHandle 也不会自动搬数据。
同进程走 UVA 和 peer access；跨进程走 cudaIpc 或 VMM POSIX FD；跨节点同一 NVLink 域走 VMM FABRIC + IMEX。
说明 P2P 的顺序、IOMMU/ACS、带宽口径，以及 SGLang Engram peer_read 与 KV Cache 分片聚合的区别。
从 FlashAttention 的 tiling 和 paged KV 看 kernel 为什么默认 K/V 在本地 HBM，以及远端 K/V 会改变哪些假设。
含 5 张手绘图，以及 CUDA Runtime / Driver API、PTX 内存模型和 SGLang 上游源码索引。

21. 手搓 Attention：从三行循环到看懂 FlashAttention

小白入门向：从三重循环、矩阵化、减 max 到分块与 online softmax，四个版本递进，每版都能跑。
数值溢出边界（fp16 分数 11.1 / fp32 分数 88.7）与四版一致性全部实测验证；含 6 张手绘图。
与第 20 篇和 TileLang FlashAttention 拆解互为对照。

参考资料

CUDA 编程简介 - 基础与实践.pdf