H100 · cuTensorMap（TMA 描述符）

目录 · ← l3 · l5 →

H100 · cuTensorMap（TMA 描述符）

本文档基于课程讲义《4. cuTensorMap.pdf》（Lesson 4 / TMA-1，41 页）整理， 系统介绍 TMA（Tensor Memory Accelerator） 与其核心抽象 cuTensorMap（描述符）。 前置阅读：H100-架构介绍.md（TMA 概述）、H100-异步与屏障.md（mbarrier、异步拷贝）。

目录

1. 什么是 TMA
2. H100 如何做到”完美异步”
3. 为什么需要描述符
4. TMA 拷贝的三步流程
5. cuTensorMap：描述符的内容
6. 创建与编码：cuTensorMapEncodeTiled
7. 数据类型与 FTZ
8. 张量秩、全局地址、维度与步长
9. boxDim 与 elementStrides
10. Shared Memory、Bank 与 Bank Conflict
11. Swizzling（地址打散）
12. Interleaving（交错布局）
13. L2 Promotion（L2 提升）
14. OOB Fill（越界填充）
15. 总结与学习衔接

1. 什么是 TMA

TMA 是 H100 新增的、用最少的线程参与实现”全异步”数据搬移的单元。

• 单个线程发起 TMA 指令后立即继续执行，整个操作由硬件在后台完成。
• 不再做地址计算，而是用描述符（descriptor）。
• TMA 能同时把数据搬到多个 SM 的 shared memory（multicast）。
• 能处理 1D–5D 张量。

2. H100 如何做到”完美异步”

• 目标：让所有单元始终忙碌。
• 手段：多个 buffer 用 TMA 加载，同时 Tensor Core 做运算（多缓冲流水线）。
• 关键点：你不需要大量复杂工程——描述符把”数据、布局、优化”全都封装好， 把重活交给硬件。

3. 为什么需要描述符

描述符是 H100 异步能力的”关键拼图”：

因为完成传输所需的全部信息都在那个 128 字节的描述符里，硬件就能在无 SM 参与的情况下后台执行。

4. TMA 拷贝的三步流程

1. 用 CUDA API 创建 cuTensorMap。
2. 用所需信息编码（encode）它。
3. 发起操作（如 cp.async.bulk.tensor）。

5. cuTensorMap：描述符的内容

cuTensorMap 是一个描述符，存储以下信息：

• 内存基指针（device 地址）
• 张量形状（每维元素数）
• 步长（字节）
• 数据类型
• 对齐与 swizzle（打散）
• 内存空间（device / host / unified）
• 阶（order）与秩（rank）（最多 32 维）
• 可选：tiling（分块）或 interleaved（交错）布局

6. 创建与编码：cuTensorMapEncodeTiled

CUtensorMap tma_desc;   // 在 host 内存创建为局部变量

• 大小 128 字节，必须 128B 对齐。
• 不是普通指针，而是 NVIDIA 定义的特定数据结构。
• 用 cuTensorMapEncodeTiled() 把值编码进 tensormap。

cuTensorMapEncodeTiled 的主要参数（本讲义展开的）：

7. 数据类型与 FTZ

dataType 决定实际从 HBM 拷贝的数据类型，TMA 引擎据此自动得到内存对齐与传输大小。

支持的类型（CU_TENSOR_MAP_DATA_TYPE_*）：

7.1 Flush-to-Zero（FTZ）

FTZ 是一种浮点优化：把非规格化数（denormal/subnormal）直接置为 0，而不是正常计算它们。

8. 张量秩、全局地址、维度与步长

8.1 Tensor Rank 与 Global Address

• Tensor Rank = 张量的维数（不是线性代数里的”矩阵秩”）。
• Global Address = 张量在 HBM 中的地址；须 16 字节对齐以获得高效访存。

8.2 Global Dimension（globalDim）

• 指定张量每维的大小（元素数，不是字节）。
• 对于 r 维数组，数据排列约定：
  • globalDim[0] = 最内层维
  • globalDim[r-1] = 最外层维
• 每个元素范围 0 到 2³²（约 40 亿）。

8.3 Global Strides（globalStrides）

• 每维之间的字节步长：硬件沿某维从一坐标移到下一坐标要跳过多少字节。
• 最内层维的步长是隐式的（由元素大小决定），因此该数组大小为 rank − 1。
• 对应关系：
  • globalStrides[0] = 第二内层维（globalDim[1]）的字节步长；
  • globalStrides[rank-2] = 最外层维（globalDim[rank-1]）的字节步长。

9. boxDim 与 elementStrides

9.1 Element Strides（elementStrides）

• 做异步拷贝时，每维想跳过的元素数。
• 大小 = rank 的数组，要求：
  • 必须非零；
  • 步长值 ≤ 8；
  • 数组大小 = rank。
• 以元素数计，不是字节。

9.2 Box Dim（boxDim）

• tile 大小：大小 = rank 的数组，每维一个条目。
• 指定每次 TMA 操作从 global memory 搬到 shared memory 的遍历盒子（traversal box）的大小。
• 以元素数计（对应所传输的数据类型）。
• 最内层维必须 ≤ swizzle 大小。

10. Shared Memory、Bank 与 Bank Conflict

• Shared memory：挂在 SM 上的片上暂存区；每个 block 分到一块，块内所有线程可读写其中任意地址，程序上是一个平坦地址空间。
• 硬件底层把它分成 32 个 bank——每个 bank 是可独立服务的”lane”。为什么是 32？因为一个 warp 有 32 个线程，理想情况是一线程一 bank。
• 关键：bank 不是 32 个手动索引的数组——你仍只算一个地址，bank 由地址决定。

10.1 Bank Conflict

• 32 个 bank 可一个 cycle 同时访问。一个 warp 发一条逻辑 load/store，但 SMEM 子系统可能分多轮执行：
  • 第 1 轮：每个 bank 各取一部分请求；第 2 轮：剩余的冲突请求……以此类推。
• 多个线程访问同一 bank → 访问串行化 → bank conflict。常见的跨步访问就会造成串行化。
• 冲突程度：1 请求 = 无冲突，2 = 2-way，4 = 4-way，8 = 8-way，32 = 最坏。

11. Swizzling（地址打散）

11.1 原理

Swizzling 是地址映射函数：把逻辑地址（程序以为写到的 GmemAddress）”打散”位， 生成物理地址（数据真正落地的 SmemAddress）。

目的：把顺序访问模式分散到所有 bank，避免冲突。

三种模式：32B、64B、128B——它们定义打散的跨度/chunk 大小，从而决定内存事务大小。 选对布局是按应用访存模式做性能调优的关键一步。

11.2 Bank 方程

bank = (a' / 4) % 32    其中 a' 是硬件实际用的 shared-memory 字节地址

• 即 bank ID 来自地址位 a'[6:2]；每 32 × 4 = 128 字节重复一次。

11.3 Swizzle 公式

a' = a ^ ((a & Y_mask) >> 3)

• a = 字节地址，Y_mask = 1 << 7，低位 4 位 a[3:0] 从不改变。
• 三种模式的位变换：
  • 32B：a[4]' = a[4] ^ a[7]
  • 64B：a[5:4]' = a[5:4] ^ a[8:7]
  • 128B：a[6:4]' = a[6:4] ^ a[9:7]

被 XOR 的位是”shared memory 中 swizzle 模式行索引的低位”，本质上不一定是数学矩阵的行； 只有当你的布局把矩阵行映射到那些 shared-memory 模式行时，它们才等于矩阵行位。

11.4 Swizzle Span 与 Atom

对 SM90，完整 swizzle 模式的重复周期：

• 每行 128 字节，拆成 8 个 16 字节单元；swizzle 只置换这 8 个单元（从一行到下一行）。
• swizzle row ≠ 你的矩阵行——它是 128 字节的 shared memory 行。
• swizzle 只以 16B 单元为粒度置换，单元内部的小段连续数据保持原样，不随机打乱。

11.5 三种模式的细节

12. Interleaving（交错布局）

12.1 为什么需要

• 全局内存里的数据不总是线性的。cuDNN 等库主要为卷积等操作使用 NCHW 内存布局。
• CUtensorMapInterleave 参数告诉 TMA 如何解码地址空间：把 HBM 中物理上交错排列的数据， 在 shared memory 里重建成逻辑上线性的布局。

12.2 两种模式（针对 NC/xHWCx 布局）

若通道数不是 slice 的整数倍，最后一个 slice 需零填充以维持交错粒度。

12.3 Interleaving 与 Swizzling 的耦合

• 二者不独立。文档规定：

  当 interleave = CU_TENSOR_MAP_INTERLEAVE_32B 时，swizzle 必须设为 CU_TENSOR_MAP_SWIZZLE_32B。

• 原因：从 global memory 解交错的 32B 块，直接喂给 shared memory 的 32B atom swizzle 逻辑， 二者步调一致地把复杂全局布局映射成无 bank conflict 的 shared 布局。

12.4 关键注意点

• 用 32B interleaving → 全局地址应 32B 对齐；用 16B → 16B 对齐。
• 用 16B interleaving → 全局步长应为 16 的倍数；用 32B → 32 的倍数。
• 用 interleaving 时，维度必须 ≥ 3。

13. L2 Promotion（L2 提升）

• 从 HBM 取数远慢于从 L2 取数。预取（prefetch） 就是发非阻塞的 HBM→L2 传输，提前准备未来数据。
• 不做提升（NONE）：内存控制器用默认 cache line 抓取（通常 32 字节）——但 L2 实际按 128 字节行管理数据，故低效。
• 把抓取尺寸提升到 128B 或 256B，一次 TMA 请求就能一次抓进更大块的连续数据，最大化带宽效率， 确保计算单元需要时数据已在 L2。

13.1 各档位与硬件含义

13.2 经验法则

• 密集 GEMM/Conv：总是用 L2_128B 或 L2_256B，最大化总线效率。
• Embedding 查表 / 稀疏：用 NONE（32B），不抓你不会碰的数据。
• 调优变量：取决于 Tensor_Inner_Dim_Bytes——若内层维 < 64B，L2_128B 可能浪费；让提升尺寸匹配你的连续数据宽度。

14. OOB Fill（越界填充）

• oobFill 处理 tensor 拷贝时的越界情况：
  • 自动在目标（shared memory）把越界区域填成 0 或 NaN，不是在 HBM 里填。
  • 源数据在 HBM 中全程不变。
  • 免去 kernel 里的手工边界检查，提升代码清晰度与性能。
• 对分块操作尤其有用（tile 超出张量边缘时）。
• 填充选择取决于需求：zero 用于 padding，NaN 用于调试。

CU_TENSOR_MAP_FLOAT_OOB_FILL_NONE;
CU_TENSOR_MAP_FLOAT_OOB_FILL_NAN_REQUEST_ZERO_FMA;  // 用 NaN 填充，但 FMA 运算时按 0 处理

第二条用 NaN 填充，但在 FMA 运算时它被当作 0。

15. 总结与学习衔接

15.1 核心脉络速记

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

15.3 一句话记忆

cuTensorMap = 一张”传输说明书”（128 字节描述符），把”数据在哪、什么类型、什么形状、怎么打散、 怎么交错、L2 怎么预取、越界怎么填”全部编码进去；TMA 读它、在后台搬数，SM 专注算矩阵。

参考来源：4. cuTensorMap.pdf（Lesson 4 / TMA-1，41 页）。