H100 · cp.async.bulk（异步批量拷贝指令族）

目录 · ← l4 · l6 →

H100 · cp.async.bulk（异步批量拷贝指令族）

本文档基于课程讲义《5. cp.async.bulk.pdf》（Lesson 5，40 页）整理， 系统介绍 Hopper H100 的 cp.async.bulk 指令族：张量拷贝、原始拷贝、multicast、prefetch、reduce。 前置阅读：H100-架构介绍.md（TMA）、H100-cuTensorMap.md（描述符）、H100-异步与屏障.md（mbarrier）。

目录

1. cp.async.bulk 概述
2. cp.async.bulk vs cp.async（Ampere）
3. 两种布局：tensor vs raw
4. tensor 布局详解
5. 完成机制：mbarrier vs bulk_group
6. L2 缓存提示（Cache Hinting）
7. 结构化 vs 非结构化拷贝
8. Multicast（多播）
9. Multicast 与 mbarrier 的正确用法
10. 动态共享内存布局
11. Multicast 完整流程
12. 结构化拷贝的操作数
13. cp.async.bulk.prefetch
14. cp.reduce.async（归约异步）
15. 总结与学习衔接

1. cp.async.bulk 概述

cp.async.bulk 是 Hopper H100 上用于硬件加速、异步批量内存传输的一族 PTX 指令。

• 这些操作被卸载到专用硬件（TMA），与 SM 的计算流水线独立执行——数据传输在后台进行，计算可继续。
• 能高效处理大块、多维张量传输：1D–5D 张量、数百到数千字节。
• 需要 barrier 对象做协调，保证异步传输与计算之间的正确顺序。

2. cp.async.bulk vs cp.async（Ampere）

核心升级：从”每线程每 16 字节算地址” → “单线程一条指令 + 描述符 + TMA 后台搬”。

3. 两种布局：tensor vs raw

3.1 Tensor 布局（结构化）

当你有 TMA 描述符（Tensor Map）、想拷贝一个特定的多维 tile 时使用：

cp.async.bulk.tensor.<ndim>.<dst>.<src>.<barrier_type>{.cache_hint}{.multicast}
    dst_addr, tensor_map, coordinate_array, mbarrier_addr;

3.2 Raw 布局（线性 / 非结构化）

用于简单、连续的字节拷贝：

cp.async.bulk.<dst>.<src>.<barrier_type>{.cache_hint}
    dst_addr, src_addr, size, mbarrier_addr;

4. tensor 布局详解

4.1 .tensor

• 加 .tensor 表示操作张量感知——工作在多维张量数据上，而非扁平数组。
• 它开启了一堆重要选项，也是启用 cuTensorMap 的前提。

4.2 .{1d,2d,3d,4d,5d}

• 告诉硬件：需要从你这里读多少个索引（数）来定位 tile。
• 一旦拿到定位 tile 所需的 n 个索引，就按 cuTensorMap 里的信息抓取 tile。
• Nd = 张量有 N 维；最多 5 维。

4.3 源/目标状态空间

cp.async.bulk.tensor.<dim>.<space1>.<space2>

• space1 = 源张量的状态空间；space2 = 目标张量的状态空间。
• 可为：global、shared::cta、shared::cluster 等。

4.4 Load Mode（.tile / .im2col）

• 这个修饰符至关重要：告诉 TMA 硬件如何解释你给的坐标、如何即时变换数据。
• .tile：TMA 用 tensorCoords 提供的坐标算出基地址，按 tensorMap 中的步长抓取一个稠密连续的多维盒子（tile）。
• .im2col：抓取时硬件加速完成 im2col 变换。
  • 过去要写 kernel 把像素从图像布局（N,C,H,W）拷成列布局（矩阵），浪费带宽与寄存器。
  • 现在只需给卷积窗口的左上角坐标，TMA 就抓取滤波所需像素、展开它们、像矩阵的一列那样放进 L2 → 加速卷积。

4.5 Completion Mechanism（完成机制）

cp.async.bulk.tensor.{}.{}.{}.<completion_mechanism>

两种完成机制：

• mbarrier::complete_tx::bytes：见 §5.1。
• .bulk_group：见 §5.2。

5. 完成机制：mbarrier vs bulk_group

5.1 mbarrier::complete_tx::bytes

• cp.async.bulk 调用带上 mbarrier 句柄（通常在 [mbar] 参数里）与完成机制说明。
• 硬件跟踪该操作，随数据搬移自动更新 barrier 的 tx-count。
• 当 tx-count 归零，barrier phase 翻转、等待的线程被释放。
• 操作数需要：mbarrier 指针 + 拷贝操作的大小。

5.2 .bulk_group

• 比 mbarrier 更简单、更轻量的替代方案。
• 不再跟踪 tx_count，而是：用 bulk_group 发起一堆拷贝 → commit_group 批量打包 → wait_group<N> 等到”最多 N 个（或更少）最近的 bulk 异步组仍 pending”。

6. L2 缓存提示（Cache Hinting）

cp.async.bulk.tensor.1d.shared::cta.global.mbarrier::complete_tx::bytes.L2::cache_hint

• 异步拷贝时若让数据淹没 L2，会踢掉你反复在用的其它重要缓存数据。为防”一次性数据”挤占 L2，用 L2 hinting。

6.1 三种策略

6.2 Loads（global → shared/dsmem）的用法

• 数据按 write-back 缓存策略加载并缓存进 L2。
• 复用权重/数据（最佳）：用 evict_last（标记持久，尽量久留）。
• 流式/一次性数据：用 evict_first（用一次就扔）。

6.3 Stores（shared → global）的用法

• 结果写回 global 时通常不需要立刻再读。
• 最终输出（最佳）：用 evict_first（写进 HBM、多半不会再在本 SM 读回），最小化缓存污染——数据进 L2（为服务写）但立刻标记为首要淘汰对象，保持缓存干净给输入用。

6.4 创建缓存策略描述符（createpolicy）

createpolicy.fractional.L2::evict_last.b64  policy_reg, 1.0;
createpolicy.fractional.L2::evict_first.b64 policy_reg, 1.0;

• 生成一个 64 位缓存策略描述符，编码特定访存模式的淘汰优先级。
• dest（policy_reg）：存放描述符的 64 位寄存器。
• fraction(1.0)：策略应用到多大比例的数据。

7. 结构化 vs 非结构化拷贝

7.1 非结构化拷贝的操作数

• shared：dstmem（shared 目标地址）、srcMem（shared 源）、size（字节）、mbar（mbarrier 指针）、cache_policy（64 位策略描述符指针）。
• global：Dst（global 目标）、Src（源）、Size、Cache_policy、Mask（掩码写：指定写目标的哪些字节）。
• dsmem：dst（DSMEM 目标）、src（shared 源）、size、completion_mechanism（mbarrier）。

  关键：这里不是搬到”分给其它 block”的 shared memory，而是搬到所有 block 共享的分布式共享内存（DSMEM）池。

8. Multicast（多播）

8.1 起源

• 算数据比搬数据快，尤其从 HBM 搬数据既耗时又耗能。
• 深度学习大量用 GEMM，而 GEMM 里很多不同 block 都要读同一块 Matrix A，去和各自的 Matrix B 块相乘。 例如 SM0、SM1、SM2、SM3 都需要 “Tile A0”。
• 旧架构：4 个 SM 各自从 global memory 请求同一份数据（如权重矩阵）——非常昂贵。
• 为什么不从 HBM 只取一次，流经总线时同时拷给 4 个 SM？ → 这就是 multicast。

8.2 操作拆解

1. TMA 从 global memory（srcMem）读 size 字节进 L2 cache；cache hint 让 L2 持久化该行，优化后续 wave 的带宽。
2. L2 控制器读一次数据，通过 cluster crossbar 广播，同时写入 mask 里每个 block 的 SMEM bank。
3. 完成时，TMA 用多播编码的 mbar 指针，原子地给每个参与 block 的 mbarrier 事务计数加 size 字节。
4. 单个 leader 线程发起这条非阻塞指令；TMA 硬件独立管理整个”取数-广播-发信号”流水线，所有 block 的线程可边等边算或休眠。

9. Multicast 与 mbarrier 的正确用法

这里很容易踩坑，要理解 cluster 内特定线程块使用 mbarrier 的规则：

• mbarrier 对象在每个参与 CTA 的 shared memory 中，以相同的相对偏移被复制。
• producer 发起 TMA 指令时，硬件把数据广播给 ctaMask 里的所有 CTA，并自动给每个目标 CTA 中那个特定地址的 mbarrier发信号。
• 硬件从指令的 mask 立刻知道组成员；接收方 CTA 不需要”arrive”来组队，只需在本地 mbarrier 实例上等待数据落地。

9.1 参与方如何协作

1. cluster 内所有 block 的第一个线程调用 expect_tx（带上它们期望的数据量），各自指向自己 shared memory 里的 mbarrier；cp.async 通过偏移找到每个 block 的 barrier。
2. 整个 block 的单个线程调用带 multicast + mask 的 cp.async.bulk，指向自己的 barrier。
3. 想只传给 block 0、3 而不传 1、2 → block 1、2 不要调用 mbarrier，也不要放进 mask。
4. 所有 block 的 consumer warp 都在本地 barrier 上 mbarrier.try_wait 自旋。
5. 到达数由 TMA 自己管理（对 mask 内的 block）：TMA 自动向指定偏移的 mbarrier 发”远程到达（remote arrival）”信号。

10. 动态共享内存布局

用动态共享内存时，要手动管理布局：

// 不要：extern __shared__ char smem[]; 然后直接往后面塞 mbarrier
// 要：
uint64_t* bar_ptr = reinterpret_cast<uint64_t*>(smem);
int tma_alignment = 128;
int data_offset = (sizeof(uint64_t) + tma_alignment - 1) & ~(tma_alignment - 1);
half* tile_ptr = reinterpret_cast<half*>(smem + data_offset);

即：先放 mbarrier（64 位），再按 128 字节对齐计算数据区偏移，避免 mbarrier 与 tile 数据重叠/错位。

11. Multicast 完整流程

以 16 个 CTA（在 cluster 内 16 个不同 SM 上）为例：

1. 建立 multicast 组：每个参与 CTA 在 shared memory 分配相同的 mbarrier 对象，调用 mbarrier.init.shared.b64（带到达数），然后对该 mbarrier 执行 arrive()。
   • 这个 arrive 不只是同步，而是硬件注册——内存子系统现在知道这 16 个 CTA 组成一个”接收相同数据”的逻辑组。
2. 每个 CTA 创建完全相同的 CUtensorMap 描述符：host 上调用 make_tma_copy(SM90_TMA_LOAD_MULTICAST{}, gmem_tensor, smem_layout, cluster_size)， 编码张量几何、数据类型、swizzle 模式，以及关键的 cluster 维度。
   • 描述符传给 kernel（标 __grid_constant__），告诉 TMA 取哪块 global、放到每个 CTA shared memory 哪里。
   • 所有 16 个 CTA 必须用完全相同的描述符——任何偏差都破坏”要同一份数据”的契约。
3. 每个 CTA（通常由单个被选中的线程）发 TMA multicast 指令：

   cp.async.bulk.tensor.shared.cluster.global.mbarrier.multicast
   

   • 注意 .cluster 作用域与 .multicast 限定符——它们向硬件表达意图。
   • 操作数：shared 目标地址、tensorMap 指针、张量坐标、mbarrier 指针，以及关键的 ctaMask。
   • ctaMask 是 16 位位掩码（对 size=16 的 cluster），bit i 表示 CTA i 是否参与；0xFFFF = 全部 16 个 CTA 接收。
4. L2 层的魔法：L2 控制器收到 16 个”看似独立”的、对同一 tile 的请求，都带相同的 mbarrier group ID。硬件识别出该模式，提升其中一个请求为 leader。
   • leader 触发一次 HBM 读（比如 1MB 权重）；数据流入 L2 时，控制器不只发给一个 SM，而是多播给全部 16 个 SM的 L1 并直接进它们的 shared memory。
   • 你付 1MB 的 HBM 带宽，却向 SM 交付了 16MB 的数据。
5. TMA 硬件随数据到达，自动递减每个 CTA 的 mbarrier 事务字节数（tx-count），跟踪完成进度。
6. 每个 CTA 执行 wait_barrier(tma_load_mbar, phase)（或等价的 mbarrier.try_wait），阻塞直到 tx-count 归零——即所有期望字节都已送达该 CTA 的 shared memory。
7. 所有 CTA 越过 barrier 后，保证 shared memory 已填好数据，可开始计算。

11.1 Multicast 指令操作数

Dst（目标地址）、Src（global 指针）、Size（操作大小）、Mbar（mbarrier 指针）、 Ctamask（16 位多播掩码）、Cache-policy（缓存策略）。

12. 结构化拷贝的操作数

12.1 global → CTA（shared）

dstMem（shared 指针）、tensorMap, tensorCoord（tensorMap 地址 + box 坐标数组）、 srcMem（global 地址指针）、cache-policy。

12.2 global → DSMEM（多播）

dstMem（DSMEM 指针）、tensorMap, tensorCoord、mbar（mbarrier 指针）、 ctaMask（选择要拷贝的 block 的掩码）、cache-policy。

12.3 TMA Stores（用 bulk_group）

tensorMap, tensorCoords（cuTensorMap 的 64 位指针 + box 坐标数组）、 srcMem（数据来源）、Cache-policy。

13. cp.async.bulk.prefetch

• 可预取数据到 L2 以降低延迟：用 cp.async.bulk.prefetch.tensor。
• 几个要点：
  • cp.async.bulk.prefetch 是性能提示：若发完 prefetch 立刻发 cp.async.bulk.tensor 而数据还没缓存， 设备仍会从 HBM 取数据。
  • 即使 main 拷贝与 prefetch 用同一个 tensorMap，L2 promotion 与 swizzling 对”数据如何缓存在 L2”没有影响。

14. cp.reduce.async（归约异步）

• 把”整个数据 tile 的原子累加”从 SM 卸载到 TMA。
• 两个版本：

  cp.reduce.async.bulk.dst.src.<completion_mechanism>{.level::cache_hint}.<redOp>.<type>
      [dstMem], [srcMem], size{, cache-policy}
  
  cp.reduce.async.bulk.tensor.dim.dst.src.<redOp>{.load_mode}.<completion_mechanism>{.level::cache_hint}
      [tensorMap, tensorCoords], [srcMem]{, cache-policy}
  

14.1 允许的 redOp 与数据类型

15. 总结与学习衔接

15.1 核心脉络速记

15.2 与本课程其他内容的衔接

15.3 一句话记忆

cp.async.bulk = 让 TMA 在后台搬数的一族指令：用 cuTensorMap 描述”搬什么、搬到哪、怎么打散”， 用 mbarrier/bulk_group 同步”搬完了没”，用 multicast 一份数据喂多个 SM， 用 prefetch 提前暖 L2、用 reduce.async 把归约也卸载给 TMA——把 SM 从”搬数”里彻底解放出来。

参考来源：5. cp.async.bulk.pdf（Lesson 5，40 页）。