学习资料
2026 夏季 InfiniTensor 训练营 CUDA 编程 3:《内存模型与规约优化》

7 月 23 日,2026 夏季 InfiniTensor 训练营 CUDA 编程方向第三课《内存模型与规约优化》开讲。
前置课程《性能模型与逐元素优化》讲了 CUDA 并行编程中的“内存墙”与 Roofline 模型,演示了向量化访存、内存对齐与合并机制等内容。

本节课承接上节并行加法性能优化内容,从 Roofline 模型出发,深入探讨如何通过提高计算强度突破性能瓶颈。重点讲解并行规约中的原子操作陷阱、分治思想应用、GPU 内存层级架构,以及 Shared Memory Bank Conflict 等底层优化细节。

回顾与性能瓶颈分析

  1. 上节课优化路径回顾
    • CPU → GPU FP32 → 访存向量化(Fold4)→ 半精度量化
    • 数据量对性能表现的影响(小数据量 vs 大数据量)
  2. Roofline 模型指导优化
    • 目标:右移 Roofline 曲线,提高计算强度 $AI = W / Q$
    • 策略:增加计算量 $W$(如算子融合)、减少访存量 $Q$

  1. 从向量加法到累加求和

  • 尴尬并行(Embarrassingly Parallel)vs 数据竞争
  • 如何解决这个竞争的问题?

原子操作与初步规约优化

原子操作(Atomic Operations)

  • 概念:不可分割的最小执行单元(All-or-Nothing)
  • CUDA API:atomicAdd 的使用与正确性验证
  • 性能代价:NCU 分析显示指令发射严重阻塞

分治规约

  • 核心思路:将全局归约分解为子问题,减少原子操作次数

1. warp 级规约

  • Lane ID 获取
  • 实现:每个 warp 内单线程循环累加 + warp 间原子操作
  • 效果:相比全线程原子操作,性能提升约 30倍

2. block 级规约的尝试与反思

  • 实现:每个 block 内单线程循环累加
  • 问题:活跃 warp 数下降,计算/内存利用率降低
  • 结论:单纯减少原子操作数不等于最优,需平衡并行度

GPU 内存层级架构(Memory Hierarchy)

内存墙问题与层级解决方案

  • 速度 vs 容量 vs 成本的权衡
  • 局部性原理:常用数据放上层,大规模数据放下层

GPU 内存结构

1. 物理硬件单元

2. 逻辑内存模型

高效规约算法演进

基于 Shared Memory 的树状规约

  • 动态共享内存分配:extern __shared__ + Kernel 启动参数配置
  • 块内同步:__syncthreads() 的作用与必要性
  • 性能:比 warp 级规约快 5.8倍,比 block 级规约快 6.4倍

Warp 规约优化

  • __shfl_down_sync:寄存器级直接通信,无需 Shared Memory 和同步
  • 优势:更高计算强度(AI 提升至 1.5),更低延迟

混合规约策略(Warp Shuffle + Tree Reduce)

  • 步骤:Warp 内 Shuffle 规约 → 结果写入 Shared Memory → 块内树状规约
  • 编译器提示:#pragma unroll 循环展开
  • 性能进一步提升

跨块规约优化

1. Two-Pass 方法:

  • Kernel1 块内规约写临时数组 → Kernel2 块间规约
  • 缺点:额外 Kernel 启动开销、隐式同步、额外访存

2. Cooperative Groups(CUDA 9.0+)

  • grid.sync():跨 Block 同步
  • 启动要求:

  • 优势:可以减少多 Kernel 调度开销,在满足硬件资源约束且同步成本较低的场景下具有性能优势
  • 注意:死锁风险、硬件/驱动版本限制

3. Two-Pass vs. Grid Sync

Shared Memory Bank Conflict 深度解析

Bank Conflict 原理

  • 在当前主流 NVIDIA GPU 架构中,共享内存划分为 32 Banks,每 Bank 4 Bytes
  • 冲突条件:同一 Warp 内多线程访问 同一 Bank 的不同地址 → 会增加共享内存访问 transaction 数量,使访问被拆分执行,导致吞吐下降
  • 非冲突特例:同一 Bank 同一地址 → 广播(Broadcast)

树状规约分析

NCU 诊断 Bank Conflict

  • Source 页面指标
  • Memory Workload Analysis

Padding 优化技巧(预告)

  • 通过内存填充改变 Bank 映射关系
  • 案例:矩阵转置中 Padding 带来 2倍 性能提升(下节课详解)

加入训练营

📺完整课程内容,请观看直播或查看课程回放:

InfiniTensor 官网https://www.infinitensor.com

B站直播: InfiniTensor 官方直播间

视频号直播: InfiniTensor 视频号预约

答疑交流: 训练营官方社群⬇️

关注与交流