GPU 高性能编程与内核优化 (High Performance GPU Programming & Kernels)
GPU 高性能编程与内核优化 (High Performance GPU Programming & Kernels)
1. GPU 架构复习 (GPU Architecture Review)
1.1 硬件层级 (Hardware Hierarchy)
- Streaming Multiprocessors (SMs): GPU 的核心计算单元。A100/H100 包含大量的 SM。
- Memory Hierarchy:
- Global Memory (DRAM): 容量大但速度慢。
- L2/L1 Caches: 速度较快。
- Register File: 每个 Thread 独有的极速存储,编写高性能代码时需大量使用。
1.2 执行模型 (Execution Model)
- Grid & Blocks:
- 计算任务被划分为多个 Thread Blocks。
- 一个 Block 被调度到一个 SM 上执行。
- Thread Blocks 之间通信昂贵(通过 Global Memory),但在同一 Block 内的 Threads 可以通过 Shared Memory 快速通信。
- Threads & Warps:
- Threads 是执行计算的最小单位。
- Warp: 32 个 Threads 组成一组,同时执行(SIMT架构)。为了性能,应尽量保证 Waves(即 Warps 的执行波次)负载均衡,。
1.3 核心性能指标 (Key Metric)
- Arithmetic Intensity: 计算量与内存访问量的比值 (Flops / Bytes)。
- Compute Bound vs. Memory Bound:
- Matrix Multiplication (矩阵乘法) 通常是 Compute Bound。
- 其他操作(如 Activation functions, Normalization)通常是 Memory Bound。
- 目标:提高 Arithmetic Intensity,减少内存搬运,。
2. 基准测试 (Benchmarking)
2.1 目的
- 测量代码的端到端 Wall clock time。
- 比较不同实现(PyTorch vs. Triton vs. CUDA)的性能差异。
2.2 关键步骤与陷阱 (Pitfalls)
- Warm-up (预热):
- 首次运行 PyTorch 代码时会涉及编译和初始化,需运行几次预热以测量稳定状态下的速度,。
- Synchronization (同步):
- CPU-GPU Asynchrony: CPU 通常会跑在 GPU 前面,将内核发射(Kernel Launch)到队列中。
- 解决方案: 必须调用
torch.cuda.synchronize()确保 GPU 完成所有任务后再停止计时,否则测量的只是 CPU 将任务推入队列的时间,。
2.3 案例分析 (Example)
- MLP Scaling: 增加层数 (Layers) 或步数 (Steps) 时,运行时间呈线性增长 (Linear scaling),。
- Matrix Multiplication: 只有矩阵足够大时才能观察到预期的性能扩展,小矩阵受限于启动开销 (Overhead)。
3. 性能分析 (Profiling)
3.1 为什么要进行 Profiling?
- Benchmarking 只能告诉你代码慢,但 Profiling 能告诉你时间花在哪里。
- 揭示 PyTorch 抽象层之下的底层 CUDA Kernels 调用情况。
3.2 工具 (Tools)
- PyTorch Profiler: 易于使用,显示高级别的 Kernel 耗时。
- NVIDIA Nsight Systems: 专业的系统级分析工具,能展示 CPU 与 GPU 的时间线交互。
3.3 关键发现 (Key Insights)
- CPU Queueing: CPU 往往比 GPU 跑得快,它会提前将多个 Kernel 放入 GPU 的执行队列中。
- Synchronization Overhead: 某些操作(如打印
loss.item())会强制 CPU 等待 GPU 计算结果,破坏流水线并行,导致性能下降,。 - Kernel Detail: 对于简单的操作(如
Add),可能会看到大部分时间花在 Kernel Launch 和数据传输上,而非计算本身。对于复杂的Matrix Multiply,底层可能会调用 cuBLAS 或 CUTLASS 的特定 Kernel。
4. 内核融合与实现 (Kernel Fusion & Implementation)
4.1 动机:Kernel Fusion
- 问题: 逐个执行操作(如 )会导致多次读写 Global Memory。
- 解决方案: 将多个算子融合为一个 Kernel(Kernel Fusion),数据只从内存读取一次,在寄存器中完成所有计算,再写回。
4.2 案例:GELU (Gaussian Error Linear Unit)
我们需要计算:
方法 A: Manual PyTorch (无优化)
- 代码: 直接使用 Python 算术运算符。
- 性能: 极慢 (e.g., ~8.1 ms)。
- 原因: 启动了大量的独立 Kernels,内存读写频繁,。
方法 B: CUDA C++ Kernel
- 实现:
- 编写
.cu文件。 - 手动计算索引:
idx = blockIdx.x * blockDim.x + threadIdx.x。 - 手动进行边界检查 (Boundary checks)。
- 通过 Python 的 C++ 扩展加载。
- 编写
- 性能: 显著提升 (e.g., ~1.8 ms)。
- 特点: 极高的控制力,但编写繁琐,容易出错(如指针越界),。
方法 C: Triton Kernel (推荐)
- 概述: OpenAI 开发的 DSL,介于 Python 和 CUDA 之间。
- 编程模型: Block-centric (以块为中心),而非 CUDA 的 Thread-centric。
- 优势:
- 自动处理 Memory Coalescing (内存合并访问)。
- 自动管理 Shared Memory。
- Python 语法,易于调试,。
- 实现细节:
- 定义
offsets为一个向量。 - 使用
tl.load和tl.store对整个 Block 进行读写。 - 数学运算直接作用于 Block 数据-。
- 定义
- 性能: 与手写 CUDA 相当 (~1.8 ms),但代码量少得多。
方法 D: Torch Compile (torch.compile)
- 原理: JIT (Just-In-Time) 编译器,自动分析 PyTorch 图并生成优化的 Triton Kernel。
- 性能: 非常快 (~1.47 ms),甚至超过了未深度优化的手写 Triton 代码。
- 结论: 对于标准操作的融合(Element-wise operations),
torch.compile通常是首选。对于复杂的、非标准的注意力机制(如 Flash Attention 3),可能仍需手写 Triton/CUDA,。
5. 进阶案例:Softmax 优化 (Softmax Optimization)
5.1 挑战
- 不同于 Element-wise 操作,Softmax 需要 Reduction (归约) 操作(求 Max 和 Sum)。
5.2 Triton 实现策略
- Row-wise Handling: 让一个 Thread Block 处理矩阵的一整行。
- Block Size: 设置为大于列数的最小 2 的幂次 (
next_power_of_2)。 - 流程:
- 加载整行数据到 SRAM。
- 计算该行的 Max。
- 计算 Exponentials 并求 Sum。
- 归一化并写回-。
5.3 性能对比
- Manual Softmax: 灾难性的慢,大量中间内存操作。
- Triton/Compiled Softmax: 单个 Fused Kernel,性能大幅提升。
6. 总结 (Summary)
- Benchmarking & Profiling: 是高性能编程的基础。永远不要凭直觉优化,要看 Profiler 数据 (Nsight Systems),。
- CPU-GPU Model: 理解 CPU 的异步提交和 GPU 的执行队列对于避免性能瓶颈至关重要。
- Kernel Writing:
- 简单融合:首选
torch.compile。 - 自定义复杂逻辑:首选 Triton (开发效率与性能的最佳平衡)。
- 极致优化/硬件特性:使用 CUDA C++,。
- 简单融合:首选
- Optimization Goal: 核心目标通常是减少 Global Memory 的访问 (Fusion),提高 Arithmetic Intensity。
注:课件中的时间数据(如 8.1ms vs 1.8ms)引用自讲座演示中的特定实验结果,实际数值取决于硬件环境。