1. 性能优化GPU 优化的核心我们为 GPU 付费的唯一原因是性能如果无法获得性能提升使用 GPU 就没有意义。理念Profile First—— 先有假设再用 profiling 工具验证。所有实验代码均在 GitHub 上运行环境推荐 Lightning AI Studio已预装 NCU 工具2. 内存层级与延迟2.1 SRAM vs DRAM内存类型位置大小延迟SRAM(Shared Memory)每个 SM 内~100 KB (可配置)~25 周期L1 Cache每 SM~128 KB~25 周期L2 CacheGPU 全局~几 MB~200 周期Global Memory (DRAM)GPU 外部~40-80 GB~290 周期发现L1 Cache 和 Shared Memory 相比 Global Memory 有~10倍的速度优势。2.2 为什么延迟难以减小文章《It’s Latency, Stupid》观点吞吐量可以通过并行化轻松提升例如 80 条电话线并行处理延迟无法通过增加并行度来降低必须从根本上改变架构GPU 的策略是**隐藏延迟hide latency**而非减少延迟通过大量线程并发执行来掩盖内存访问延迟。3. 合并全局内存访问 (Coalescing Global Memory Accesses)3.1 为什么重要GPU 是吞吐量导向架构理想情况下读取 50 个连续元素不比读取 1 个元素慢多少。但前提是内存访问必须是合并的。3.2 合并 vs 非合并// 合并版本Coalesced __global__ void copy_data_coalesced(float* out, float* in, int n) { int idx blockIdx.x * blockDim.x threadIdx.x; if (idx n) { out[idx] in[idx]; // 连续访问 in[idx] } } // 非合并版本Non-Coalesced __global__ void copy_data_non_coalesced(float* out, float* in, int n) { int idx blockIdx.x * blockDim.x threadIdx.x; if (idx n) { out[idx] in[idx * 2 % n]; // 跳跃访问间隔访问 } }3.3 性能对比 NCU profiling指标非合并版本合并版本DRAM Throughput~90%~82%L1 Cache Throughput~30%~37%Kernel Duration~764 μs~5582 μs非合并版本慢了约4倍L1 Cache 命中率从 37% 降到 30%。3.4 在 PyTorch 中的表现非合并访问在 PyTorch 中通常以stride形式出现# 物理内存连续但按 stride2 读取xtorch.randn(1024)yx[::2]# stride2非连续访问模式4. 最大化占用率 (Maximizing Occupancy)4.1 占用率占用率 实际运行的 warp 数量 / GPU SM 能支持的最大 warp 数量。目标让 SM 始终有 warp 可调度这样当某个 warp 等待内存时另一个 warp 可以执行。4.2 占用率不足的问题量化效应Tile Quantization平铺量化当矩阵维度无法被线程块大小整除时最后一行的 tile 会产生不完整的 tile导致资源利用不足。Wave Quantization波次量化当总 tile 数无法被 SM 数量整除时最后一批 tiles 会不均衡地分配到 SM 上。# 示例M1024, N1024, 改变 K 值# K1012: 耗时 4x差# K1016: 耗时 1x优4.3 最佳矩阵维度Tensor Core对于 A100 上的 Tensor Core矩阵维度最好是 16 的倍数数据类型最小维度倍数原因FP6488 bytes per elementTF32164 bytesFP16/BF16322 bytesINT8641 byte这解释了为什么 PyTorch 社区广泛使用 padding如 vocab32000 是 64 的倍数。4.4 CUDA Occupancy CalculatorCUDA 提供了便捷的工具来计算最佳 launch 参数#includecuda_occupancy.hintmain(){intblock_size;intmin_grid_size;// 传入 kernel 函数引用自动计算最优配置cudaOccupancyMaxPotentialBlockSize(min_grid_size,block_size,copy_data_coalesced,0,// 共享内存大小无额外需求0// 最大 block size0 表示无限制);printf(Recommended block size: %d\n,block_size);printf(Minimum grid size: %d\n,min_grid_size);}运行结果示例Recommended block size: 1024 Minimum grid size: 40不同 GPU 上的结果不同如 A100 推荐 block_size768, grid_size160是 T4 的 4 倍。4.5 如何诊断占用率问题查看 NCU 报告中的Theoretical Occupancy vs Measured OccupancyTheoretical occupancy: 100% Measured occupancy: 77%差距可能来源Warp 调度开销Block 间/Block 内负载不均衡5. 计算密集 vs 内存密集分析 (Roofline Model)5.1 Roofline Model 概述Performance (GFLOPS) ^ / | / | Compute / | Bound / | Region / | (平坦) / | / / | / / | / Memory / | / Bound / | / Region / | /(斜线) / | / / ------------------------- Operational Intensity (FLOPs/Byte)X 轴运算强度 总 FLOPs / 总内存访问字节数Y 轴实际性能 (GFLOPS)两个区域Memory Bound Region左侧斜线部分性能受内存带宽限制Compute Bound Region右侧平坦部分性能受 GPU 算力限制5.2 运算强度计算示例 1ReLU 激活函数# ReLU: y[i] max(x[i], 0)# 场景 A: x[i] 0 需要写入 0# 读取 1 个 float32 (4 bytes)# 写入 1 个 float32 (4 bytes)# 执行 1 个比较操作运算强度1FLOP/8bytes0.125# 场景 B: x[i] 0 无需写入最乐观情况# 读取 1 个 float32 (4 bytes)# 执行 1 个比较操作运算强度1FLOP/4bytes0.25结论运算强度 1 的操作如 ReLU是内存带宽密集型。示例 2矩阵乘法 (GEMM)# C[M×K] A[M×N] B[N×K]# M1024, N1024, K1024# 总 FLOPs M × K × 2N 1024 × 1024 × 2048 ≈ 2.1B# 总内存访问 A(M×N) B(N×K) C(M×K)# 1024² 1024² 1024² 3 × 1024² ≈ 3MB运算强度2.1B/3MB ≈700FLOPs/Byte结论大规模矩阵乘法是计算密集型但小矩阵如 1×1, 2×2会变成内存密集型。作业尝试推导矩阵向量乘法的运算强度理解为什么它通常是内存密集型。5.3 优化方向对照表瓶颈类型优化策略具体手段Memory Bound提高运算强度Kernel Fusion融合多个小操作Quantization减少数据传输量Thread Coarsening单线程做更多工作Compute Bound改进算法更好的矩阵乘法算法更高效的数值计算建议使用torch.compile可以自动完成 Kernel Fusion 和编译优化。6. 最小化控制流分歧 (Minimizing Control Divergence)6.1 问题CUDA 按warp32 个线程为单位调度指令。一个 warp 内的所有线程必须同时执行同一条指令。// 分歧示例 if (data[idx] % 2 0) { out[idx] data[idx] * 2; // 偶数乘以2 } else { out[idx] data[idx] 1; // 奇数加1 }问题某些线程执行*2其他执行1快的线程必须等待慢的线程效果是乘法的而非加法的性能损失6.2 重写消除分歧将条件分支重写为代数运算// 重写版本无条件分支 bool is_even (data[idx] % 2 1) ^ (data[idx] % 2 0); out[idx] data[idx] * is_even (data[idx] 1) * (1 - is_even); // 更简洁的写法 // 如果 is_even1: out data[idx] * 2 // 如果 is_even0: out data[idx] 16.3 性能对比版本分歧开销Duration原版有分支98,000 branch instructions0.74 ms重写无分支65,000 branch instructions0.24 ms约 3 倍加速这是最 impactful 的优化之一。6.4 注意循环导致的分歧不严重只有 warp 边界处可能分歧嵌套 if 语句很危险每层分支指数级增加分歧概率NVCC 编译器会尝试自动消除某些分歧但无法完全消除7. 线程粗化 (Thread Coarsening / Vectorization)传统观点每个线程做尽可能少的工作。新观点在内存带宽密集型场景下让单个线程做更多工作可以大幅提升性能。7.2 示例向量加法// 细粒度版本 __global__ void vector_add(float* c, float* a, float* b, int n) { int i blockIdx.x * blockDim.x threadIdx.x; if (i n) { c[i] a[i] b[i]; } } // 粗化版本Coarsening Factor 2 __global__ void vector_add_coarsened(float* c, float* a, float* b, int n) { int i (blockIdx.x * blockDim.x threadIdx.x) * 2; if (i 1 n) { c[i] a[i] b[i]; c[i 1] a[i 1] b[i 1]; } else if (i n) { c[i] a[i] b[i]; } }7.3 性能对比版本Duration细粒度~23 μs粗化 (factor2)~0.2 μs约 100 倍加速演讲现场实测约 30 倍这是本次课程中性能提升最大的优化。7.4 原理分析粗化后DRAM Throughput: 从 ~80% 降到 ~1%数据可能完全 fit 在 L2/L1 cache单线程读取更多连续数据提高内存访问效率tips实测结果受数据大小和缓存状态影响有用户测试了 factor4 和 factor8未见显著额外提升。8. 私有化 (Privatization)将部分更新的数据存储在私有副本寄存器或 shared memory中最后再写回全局内存避免频繁访问 global memory。示例// 非私有化每次操作都访问 global memory for (int i 0; i n; i) { sum input[i]; output[i] sum; // 频繁写回 } // 私有化在 shared memory 中累积最后写回 __shared__ float private_sum[SUBTILE_SIZE]; // ... 累积操作 ... // 最后一次性写回 output应用Sliding WindowSliding Window Attention 是私有化的典型应用不计算完整的 N×N 注意力矩阵只计算局部窗口如 window_size8内的注意力大幅减少内存访问这是 Mistral 和 Mixtral 等模型使用的重要技巧。9. 数学重写在线 Softmax 与 Flash Attention9.1 标准 Softmax 的问题# 标准 Softmax 算法# 第一遍计算分母所有 exp 的和totalsum(exp(x[i])foriinrange(n))# 第二遍计算输出y[i]exp(x[i])/total问题需要读取数据两次一遍算 total一遍算 y数值溢出风险exp(1000) 会溢出在低精度FP16下尤其严重9.2 安全版 Softmax防溢出mmax(x)# 减去最大值防止溢出y[i]exp(x[i]-m)/sum(exp(x[j]-m))问题现在需要读取三次max、total、exp9.3 在线 SoftmaxOnline Normalizer核心思想维持一个假的归一化因子随着新数据的到来渐进修正。# 论文Online Normalizer Calculation for Softmax# https://arxiv.org/abs/2002.09018# 变量定义M_prevpreviousmaxM_currcurrentmaxL_prevprevioussumof exp L_currcurrentsumof exp# 关键公式L_curr(L_prev*exp(M_prev-M_curr))exp(x_curr-M_curr)核心洞察如果 max 没变新数据直接加到 sum 上如果 max 变了用exp(M_old - M_new)来 rescale 旧的 sum9.4 效果版本内存访问次数标准 Softmax2 reads 1 write安全版 Softmax3 reads 1 write在线 Softmax2 reads 1 write无溢出风险9.5 与 Flash Attention 的关系Flash Attention 使用了在线 Softmax 的思想将 Q、K、V 分 tile 加载到 shared memory使用在线归一化器逐步计算 softmax无需完整保存 N×N 注意力矩阵节省 O(N²) 显存10. TilingJeremy 在 Lecture 2 中已详细讲解本文仅做简要回顾在矩阵乘法中某些元素会被多次复用矩阵 A 的元素在计算 C 的一整行时会被用到 K 次矩阵 B 的元素在计算 C 的一整列时会被用到 M 次优化策略将热点数据加载到 shared memory减少 global memory 访问与私有化的关系Tiling 是私有化的一种特殊形式因为其重要性而独立讨论。实现要点Tiling 算法本质是 4 层嵌套循环外层两个矩阵维度 × 内层两个 tile 维度实现并不复杂。11. 综合优化技巧受益类型说明Coalescing Memory AccessMemory Bound最基础、最常见Maximizing Occupancy两者皆可使用 occupancy calculatorThread CoarseningMemory Bound本课程最大加速来源Minimizing Divergence两者皆可消除分支分歧是乘法效应Tiling / Data ReuseMemory Bound减少全局内存带宽PrivatizationMemory Bound减少 global memory 访问Math Rewrite两者皆可利用数值恒等式QuantizationMemory Bound降低内存带宽需求12. 工具与资源12.1 Profiling 工具工具用途NCU(NVIDIA Nsight Compute)Kernel-level profilingNVTX代码中添加性能注释标记PyTorch Profiler上层应用分析12.2 论文Citadel GPU Benchmark Paper- 理解 GPU 微架构Demystifying Nvidia Ampere Architecture- 微基准测试方法Online Normalizer Calculation for Softmax- 数学重写示例Flash Attention Paper- 在线归一化在注意力机制中的应用Programming Massively Parallel Processors(Ch. 6) - 所有优化技巧的来源12.3 讲座Bill Dally (NVIDIA Chief Scientist) 的所有 YouTube 讲座 - 理解 GPU 硬件设计哲学13. 总结CUDA 优化的三个核心问题瓶颈在哪里使用 NCU 分析确定是Compute Bound还是Memory Bound为什么在这里使用 Roofline Model理解运算强度我能做什么对照优化清单选择合适的技巧优秀的工程师兼具数学直觉和系统能力这也是 CUDA Mode 社区推崇的方向。