关卡阅读材料 // 01
内存层次
线程已经足够多之后，下一步不是继续增加线程，而是让一个 Warp 的内存请求更集中，并把会重复使用的数据放到合适的位置。
建议边读边写下答案；时长包含代码推演与练习。
进入互动关卡→- 01导读与目标05 MIN
- 02心智模型08 MIN
- 03概念深挖15 MIN
- 04代码推演12 MIN
- 05历史生态08 MIN
- 06练习复盘12 MIN
不仅让程序在 GPU 上运行，还要使用行业工具识别内存瓶颈，以实验数据推动优化。
关卡把合并访问、数据复用和 Shared Memory 变成可见的数据路径；实验要求用 Profiler 证明事务、吞吐与 Bank Conflict 是否真的改善。
- 完成课程 00 或能编写基础 Kernel
- 理解 Warp 与全局索引
- 会比较参考输出
- Nsight Compute
- CUDA C++
- CUDA C++ Best Practices Guide
导读与目标
先记住这一句
GPU 性能不仅取决于算多少，还取决于数据以什么形状到达计算单元。
完成后你应该能够
- 01
从一个 Warp 的地址集合判断访问是否可能产生过多全局内存事务。
- 02
根据数据复用范围选择寄存器、Shared Memory 或全局内存。
- 03
解释 Halo、协作加载与同步在邻域计算中的角色。
- 04
识别 Bank Conflict、错误同步与过度缓存等常见问题。
建立心智模型
- 01地址
一个 Warp 同时发出 32 个地址；硬件会把它们合并成尽可能少的内存事务。
- 02复用
若同一 Block 多次使用相邻数据，可先协作加载到 Shared Memory，再从片上复用。
- 03同步
只有确认所有参与线程完成加载后，其他线程才能安全读取完整 Tile。
概念深挖
合并访问
相邻线程访问相邻、适当对齐的数据时，通常能减少为服务一个 Warp 所需的全局内存事务。
作用域决定位置
寄存器属于线程，Shared Memory 属于 Block，全局内存可被整个 Grid 与 Host 访问。数据共享范围决定存放位置。
Shared Memory 是程序管理的缓存
它很快但容量有限，也可能出现 Bank Conflict；需要显式加载、同步与布局设计。
从地址数量转向事务数量
一个 Warp 执行一次全局内存 load 时，会同时提出一组地址。硬件按所需的内存段组织事务，因此性能问题不是简单的“读取了 32 个值”，而是“为了得到这些值搬运了多少段数据”。连续地址通常能提高有效字节与传输字节之比。
合并规则会随 Compute Capability、数据宽度和缓存路径变化，所以不要把某个固定事务数字当成永久规律。稳定的推理方式是观察同一 Warp 的地址分布、对齐和实际使用的字节，并用 Nsight Compute 等工具验证。
停下来想一想32 个线程读取连续 float，和读取相隔 4 KB 的 float，哪一种通常浪费更多传输？＋
参考答案相隔 4 KB 的访问通常落到更多内存段，产生更多事务与无用传输；连续访问更容易合并。
Tile 把重复搬运变成协作
邻域、卷积与矩阵算法常让相邻线程重复读取同一输入。Shared Memory Tile 的思想是让 Block 先以合并方式加载一块数据，再让多个线程从片上位置复用。Halo 是为了让 Tile 边缘线程仍能访问邻居而额外加载的边界数据。
Tile 大小要同时考虑复用收益、Shared Memory 容量、Block 大小与边界开销。Tile 过小会让 Halo 占比变高，过大会限制驻留 Block 数。真正的优化需要比较减少的全局请求与新增同步、索引和资源成本。
停下来想一想64 线程做三点邻域，为什么 66 个输入值就可能服务 64 个输出？＋
参考答案64 个主体位置加左右两个 Halo 覆盖了所有相邻访问；主体值可以被相邻线程复用，而不必每个线程各读三次全局内存。
同步必须被整个参与组一致到达
__syncthreads() 是 Block 级屏障：只有当所有参与线程都到达后，任何线程才能继续。把它放在只有部分线程执行的分支里，可能让一部分线程等待永远不会到达的伙伴。正确模式通常是条件化加载，但无条件到达屏障。
Shared Memory 由 Bank 组织。不同线程访问映射到同一 Bank 的不同地址时，请求可能被串行化；对二维 Tile 增加 Padding 是常见处理方式。同步保证可见性，却不会自动消除 Bank Conflict 或索引错误。
停下来想一想if (threadIdx.x < 16) __syncthreads(); 为什么危险？＋
参考答案同一 Block 中只有部分线程进入屏障，其余线程不会到达，行为未定义并可能死锁。
把概念放回代码
01__shared__ float tile[BLOCK_SIZE + 2];02int i = blockIdx.x * blockDim.x + threadIdx.x;03tile[threadIdx.x + 1] = input[i];04// 边界线程再加载左右 halo05__syncthreads();06output[i] = tile[t] + tile[t + 1] + tile[t + 2];代码逐段推演
- 01ALLOC为主体与 Halo 分配 Tile
BLOCK_SIZE + 2 为左右各留一个额外位置；二维问题通常会在两个维度加入边界。
- 02LOAD让相邻线程合并读取
主体数据按 threadIdx.x 映射到连续地址；边界线程再承担少量 Halo 加载。
- 03BARRIER在消费前完成协作
所有线程都完成自己的加载职责后统一 __syncthreads()，确保 Tile 对整个 Block 可见。
- 04REUSE从片上位置组合输出
每个输入只需从全局内存加载一次，却可被多个相邻输出读取；最后仍要处理全局边界。
知识图谱
补充参考 · 不要求一次读完，也不计入核心 60 分钟
核心术语
- Memory Transaction
- 硬件为满足一个 Warp 的地址请求而搬运的内存段；程序请求的字节数不等于总传输字节数。
- Coalescing
- 让 Warp 的地址落在尽可能少且对齐的内存段中，提高有效字节占比。
- Shared Memory
- 由 Block 内线程显式管理和共享的片上存储，适合协作加载与复用。
- Halo
- Tile 边缘为计算邻域输出而额外加载的输入区域。
- Bank Conflict
- 同一请求中的线程访问同一 Bank 的不同地址，导致 Shared Memory 请求被拆分或串行化。
带数字推演
- 01直接读取
64 × 3 = 192 次标量读取请求，相邻线程会重复请求许多输入。
- 02协作 Tile
64 个主体值加左右各 1 个 Halo，共加载 66 个值到 Shared Memory。
- 03算复用
192 ÷ 66 ≈ 2.9，代码层面的请求数量约降到原来的三分之一。
诊断手册
带宽利用率低且 Load 请求很多
- 先检查
- Warp 地址是否连续、对齐，数据布局是否让相邻线程访问相邻元素
- 需要的证据
- Nsight Compute 的内存事务、吞吐与 Warp 地址模式
加入 Shared Memory 后更慢
- 先检查
- 复用次数、额外同步、Halo 比例、Shared Memory 容量与 Occupancy
- 需要的证据
- 优化前后全局字节、屏障等待和活跃 Warp 对比
二维 Tile 出现异常 Shared Memory 延迟
- 先检查
- 行跨度是否让多个线程映射到同一 Bank
- 需要的证据
- Bank Conflict 指标，以及增加一列 Padding 前后的对照实验
从硬件到生态
- 01全局内存主导
早期 GPGPU 需要开发者严密组织地址来获得带宽。
合并访问成为第一代通用 GPU 优化核心。 - 02缓存 + 可编程 Shared Memory
缓存层次、ECC 与片上数据复用能力增强。
优化从地址连续性扩展到复用、可靠性和资源交换。 - 03异步数据搬运
更现代的架构提供更强的异步复制与 Tile 搬运机制。
软件开始把“加载下一块”和“计算当前块”组成显式流水。
硬件与生态坐标
Fermi 强化了面向通用计算的缓存、ECC 与更完整的内存系统。它让“线程如何取数、缓存如何工作、可靠性如何保证”成为数据中心 GPU 编程的重要组成，而不是只看算术单元数量。
练习与复盘
一个 Warp 的线程读取 input[i * 16]，最先应该检查什么？
- A. 地址是否分散到许多内存事务
- B. Block 数是否一定等于 16
- C. 是否改成更多 CPU 线程
查看答案＋
跨步访问会把同一 Warp 的地址拉散。先观察请求形状，再决定是否重排数据或用 Shared Memory。
一个 Warp 读取 32 个连续 float，但首地址偏移一个 float。这一定与对齐访问同样高效吗？
提示＋
考虑访问跨越的内存段。
参考答案＋
不一定。错位可能多覆盖一个内存段，增加事务；缓存复用有时会减轻影响，但应检查对齐并测量。
64 个线程各自从全局内存读取左、中、右三个值，共提出多少次标量读取？协作 Tile 至少加载多少个值？
提示＋
忽略缓存，把每次表达式读取都计数。
参考答案＋
直接方式为 192 次标量读取；一维 Tile 为 64 个主体加 2 个 Halo，共 66 个值。事务数量仍取决于访问与硬件。
只有负责 Halo 的线程执行 __syncthreads()。如何修复？
提示＋
把加载职责和屏障参与分开。
参考答案＋
条件分支只包住 Halo 加载；把 __syncthreads() 移到分支之后，让 Block 中所有线程一致到达。
一个数据只被当前线程使用一次，应该优先放进 Shared Memory 吗？
提示＋
看共享范围与复用。
参考答案＋
通常不应。线程私有且短生命周期的值更适合寄存器。Shared Memory 适合 Block 内跨线程共享或显著复用的数据。
可选实践实验
不计入核心 60 分钟 · 需要相应 CUDA / GPU 环境
实验目标
用证据比较连续、跨步与 Shared Memory Tile
构造三种访存版本，在保证结果一致的前提下，用 Nsight Compute 区分地址模式、全局内存效率与 Shared Memory 行为。建议步骤
- 01
实现 input[i] 与 input[(i * stride) % n] 两个读取版本，固定数据规模、Block 配置与预热次数。
- 02
用 Nsight Compute 收集全局 Load/Store、Memory Throughput 与 Warp 相关指标，并保存命令和报告。
- 03
为三点邻域增加 Shared Memory Tile 与 Halo，检查所有线程是否一致到达屏障。
- 04
比较直接读取与 Tile 版本；若优化无效，先解释复用、规模或资源占用，而不是修改结论。
提供三个版本的正确性、Profiler 指标表和一句因果解释：地址/复用发生了什么，指标如何支持这个判断。
构造一个二维转置 Tile，比较无 Padding 与 +1 Padding 的 Bank Conflict。