1. 课程导入：学习目标、课程结构、GPU内核剖析与优化的重要性

1.1 为什么需要学习GPU内核剖析

现代高性能计算与人工智能训练高度依赖GPU的并行计算能力。然而，绝大多数开发者在编写CUDA、HIP或OpenCL内核时，往往只关注“功能正确”——即计算结果是否与预期一致，却很少追问：这个内核真的把GPU的算力用足了吗？ 一个典型的例子是：同一份矩阵乘法代码，在未优化时可能仅达到GPU峰值性能的5%～15%，而经过内核剖析与针对性优化后，可以提升至60%～80%甚至更高。这种数量级的差距，正是由访存模式、线程组织、占用率、指令流水线效率、缓存命中率等底层因素决定的。

GPU内核剖析（Kernel Profiling）就是通过性能计数器、时间戳采样、指令追踪等手段，量化内核执行过程中各类硬件事件的发生频率与耗时分布，从而定位性能瓶颈的技术过程。它不同于简单的“计时跑完”，而是深入到SM（流多处理器）、内存层级、指令调度器等微观层面，回答诸如：“L2缓存命中率为何只有40%？”“全局内存合并访问是否被破坏？”“是否存在严重的线程束发散？”“SM占用率是否过低导致延迟无法隐藏？” 这类问题。

1.2 课程学习目标

本课程面向具备基础CUDA编程经验、希望系统掌握GPU性能分析方法论的开发者。完成本课程后，你将能够：

• 读懂并解释NVIDIA Nsight Compute、Nsight Systems、nvprof（旧版）等工具输出的关键性能指标，包括但不限于：SM占用率、Achieved Occupancy、Warp Stall原因分布、Memory Throughput、Compute Throughput、L1/L2命中率、DRAM带宽利用率等。
• 建立从“指标异常”到“代码原因”的映射能力：例如，看到high local memory pressure能联想到寄存器溢出或数组索引非编译期常量；看到low warp occupancy能联想到寄存器使用过多或共享内存超限。
• 掌握系统化的内核优化方法论：包括访存优化（合并访问、共享内存复用、只读缓存）、执行配置优化（block大小、grid大小、寄存器上限控制）、指令级优化（减少整数除法、使用向量化类型）等。
• 能够独立完成一次完整的性能剖析报告：从环境搭建、基准测试、瓶颈定位、优化实施到前后对比，形成可复现的优化记录。

1.3 课程结构总览

本课程共30章，按知识递进关系分为五个阶段。下表展示了整体结构及各阶段重点：

每一章均包含：原理讲解 → 工具操作演示 → 案例剖析 → 常见误区与复盘四个固定模块，确保理论与实践不脱节。建议学习时同步打开Nsight Compute，跟随案例操作，而非仅阅读文字。

1.4 GPU内核剖析的重要性：三个真实场景

为了直观说明剖析的价值，我们看三个典型场景：

场景一：科学计算中的Stencil计算（如地震波模拟）。某团队实现了一个7点Stencil内核，初始性能仅为理论峰值的8%。通过Nsight Compute剖析发现：全局内存加载的合并度仅为32%（理想应为100%），原因是数据布局为AoS（Array of Structures）而非SoA（Structure of Arrays）。仅将布局改为SoA并添加__restrict__关键字后，性能提升4.2倍。

场景二：深度学习中的LayerNorm内核。某推理框架的LayerNorm在A100上耗时120μs，剖析显示其SM占用率仅25%，原因是每个线程块使用了大量寄存器（超过128个），导致每个SM只能驻留少量线程块。通过使用__launch_bounds__(256, 8)限制寄存器使用，并改用Warp Shuffle进行跨线程归约，耗时降至45μs。

场景三：图数据库中的稀疏矩阵向量乘（SpMV）。性能剖析发现DRAM带宽利用率仅30%，而L2缓存命中率高达95%——说明数据局部性尚可，但线程负载不均衡导致部分SM空闲等待。采用行平衡调度（按非零元数量分段）后，吞吐量提升2.8倍。

这三个案例的共同点是：问题无法通过代码审查直接发现，必须依赖剖析工具提供的硬件计数器数据。没有剖析，开发者可能花费数周尝试各种“看似合理”的优化，却收效甚微。

1.5 常见误区与学习建议