GPU 高性能计算

目录

  • 计算的硬件加速思路
    • 思路1: 缩短指令的执行时间
    • 思路2: 提升单位时间内指令的吞吐量
      • 对于cpu scalar 实现
      • CPU 使用 SIMD 实现 CPU 的向量化加速：
      • Cuda SPMD
  • 并行策略
    • 指令并行 ILP (Instuction Level Parallel)
      • CPU 指令并行示例
      • GPU 指令并行示例
    • 多核并行 Multiple core parallelism
      • CPU 多核并行示例
      • GPU 多核并行示例
    • Multi-threaded parallelism
  • GPU 编程模型
    • GPU执行模型
    • GPU 内存模型
    • GPU 计算：Cuda Core VS Tensor Core
    • Cuda 编译原理
      • CUDA NVCC 工作原理及 Host/Device 编译器机制详解
      • NVCC 的编译流程 (Compilation Trajectory)
      • 第三步：Host 编译 (CPU 侧)
      • 第四步：胖二进制合并 (Fatbinary Embedding)
      • 第五步：链接 (Linking)
  • Roofline Model —— 性能优化方向的理论指导
    • 核心公式：性能的物理边界
    • 关键概念：算术强度 (Operational Intensity)
    • 图像解读：斜率与平台
    • 实践分析
      • 硬件参数获取（以 RTX 4000 Ada 为例）
      • 实战案例分析
    • Roofline model的局限性
  • 内存访问密集型高性能计算
    • GPU 的内存层次结构 (Memory Hierarchy)
      • L1 缓存的特殊之处
      • 各级存储访问延迟的直观感受
    • 内存访问合并
    • Bank conflict
      • 优化技巧
    • 内存访问密集型高性能实战
      • Naive 实现
      • 共享内存优化(Shared Memory)
      • 分析与对比
  • 高级 cuda 特性 —— 从矩阵乘法(gemm)说起
    • 使用Roofline Model对矩阵乘法进行性能分析。
      • 情况 A：无缓存/无复用的实现 (Naive)
      • 情况 B：理想状态/完全复用 (Ideal Cache)
      • 从 Memory-bound 到 Compute-bound 的转变
    • 显性的矩阵乘法性能分析
      • 基本数据
      • 理论极限分析
      • 缓存缺失与带宽约束分析
      • 综合对比总结
    • 矩阵乘法中的分块与数据复用
      • Shared Memory Tiling (块级复用)
      • Register Tiling (线程级复用)
    • 利特尔法则与占用率控制
      • 利特尔法则 (Little’s Law)：
      • 占用率 (Occupancy) 的计算
      • 人工控制占用率
      • 总结
    • 调度优化
      • 隐藏访存延迟（Hiding Latency）
      • 重叠计算与搬运（Overlapping Data Movement & Computation）
    • Tensor Cores
      • 什么是 Tensor Core？
      • TF32 (TensorFloat-32) 精度
      • 关键指令详解：mma.sync
      • 寄存器布局（Register Layout）—— 最难点
      • wgmma：warp-group level tensor core instructions
  • GPU 上动态的调度 —— 从run-length-compression说起
    • 并行扫描问题
      • Warp scan
      • Block scan
      • grid 级 scan：递归扫描 block sums
      • 整体scan 流程
      • workspace size
    • Run length compression
    • Run length decompression
  • Cuda TMA （Tensor Memory Accelerator）
    • 基础介绍
    • Tensor descriptor：CUtensorMap
    • TMA 完成异步数据搬运
      • 单 block、单 lane、单 tile 的 TMA load
      • TMA store
      • TMA reduce，把“搬运”和“归约”合并
    • TMA 中使用 swizzling pattern
  • TPU 与分布式通信原语
    • TPU 的硬件模型
    • TPU Pallas
    • TPU RDMA
    • reduce-scatter 与 all-gather
    • collective matmul：通信与计算重叠
      • 为什么需要通信与计算重叠
  • TIRx 基础
    • 为什么是 TIRx
    • 从一个最小 GEMM 开始，TIRx 在做什么
    • 理解 TIRx，关键是看三个决定
      • Layout API
      • Named Axes：Layout 不一定产生线性地址
      • Swizzle 单独建模
    • TIRx 真正重要的地方，不在于语法
  • TIRx 与高性能GEMM (上)
    • 从 GEMM 公式到数据路径
    • Step 1：单 tile GEMM，先把完整路径跑通
    • Step 2：加入 K-loop，真正的难点是 barrier phase
    • Step 3：空间 tiling，把单 CTA 扩展成二维 grid
    • Step 4：TMA 完成 GMEM 到 SMEM 的异步数据搬运
    • Step 5：多 stage pipeline 实现计算和数据搬运的重叠
    • Step 6：Persistent kernel 让 CTA 持续领取 tile
  • TIRx 与高性能GEMM (下)
    • step7: warp specialization 和流水线化
    • Step 8：两个 CTA 组成 cluster，扩大片上复用半径
    • Step 9：增加第二个 MMA consumer，让 B tile 被更多次使用
    • 性能对比
    • 总结
  • Flash Attention 4：在 Blackwell 上把 Softmax 接进 Tensor Core 流水线
    • Attention 的真实瓶颈：不是公式，而是中间矩阵
    • 在线 Softmax 与 Conditional Rescaling
    • 一个 K/V Block 的数据流
    • Warpgroup 分工：四个 WG 各管一段依赖链
    • Softmax、Correction 与 Barrier 的交接
    • Conditional Rescaling 为什么有效
    • Causal Attention 与 GQA：同一条数据路径上的两个特化
    • 调度策略：non-causal 和 causal 的任务形状不同
    • 当前实现与论文描述的差异
    • 编译与验证约束
    • 总结
  • Triton 高性能编程
    • Vector Add：从一个向量加法 kernel 看懂编程模型
      • Compute Kernel：一个 program 处理一个连续 block
      • 为什么 mask 是必需的
      • Launch Wrapper：Python 张量如何进入 GPU kernel
      • Scope / Layout / Dispatch：用三个问题读这段代码
      • Benchmark
      • 小结
    • Fused Softmax：看懂 Triton 的算子融合
      • Softmax 的数学式与内存流量
      • Compute Kernel：一个 program 处理若干行
      • Power-of-two padding：为什么尾部要填负无穷
      • Wrapper：从 shape 推导 launch 数量
      • Scope / Layout / Dispatch：和 Vector Add 相比发生了什么
      • Benchmark：用 GB/s 衡量融合后的内存效率
      • 小结
    • Triton Matrix Multiplication：从 block 映射到 L2 友好的 GEMM
      • 为什么要自己写 matmul
      • 把 GEMM 拆成 program 级 tile
      • Pointer Arithmetic：二维坐标如何变成一组地址
      • Program ID 到输出 tile 的映射
      • L2 cache reuse：为什么 grouped ordering 能提速
      • Kernel 主体：load、dot、accumulate、store
      • Autotune：把 tile shape 变成可搜索参数
      • Wrapper、正确性和 FP8 特化
      • Benchmark：结果应该怎样读
      • Scope / Layout / Dispatch 总结
  • 参考
GPU 高性能计算
  • Search

© Copyright 2026, Pin Fang, .