H100 · Clusters、数据类型、内联 PTX、状态空间

目录 · ← l1 · l3 →

H100 · Clusters、数据类型、内联 PTX、状态空间

本文档基于课程讲义《2. Clusters, Data types, inline PTX, State Spaces.pdf》（Lesson 2，47 页）整理， 覆盖四个主题：线程块簇（Thread Block Clusters）、数据类型（Data Types）、 内联 PTX（Inline PTX）、状态空间与指针（State Spaces & Pointers）。 前置阅读：H100-架构介绍.md（尤其 GPC、DSMEM、Tensor Core 相关内容）。

目录

1. 线程块簇（Thread Block Clusters）
2. 分布式共享内存（DSMEM）
3. 创建与使用线程块簇
4. PTX 与内联 PTX
5. PTX 状态空间（State Spaces）
6. 数据类型（Data Types）
7. 内存地址与 Shared Memory Bank
8. 指针与状态空间
9. 总结与学习衔接

1. 线程块簇（Thread Block Clusters）

1.1 什么是线程块簇

Cluster 是”最多 16 个线程块”的集合，保证被共同调度到同一 GPC 内相邻的 SM 上并发执行。

它把传统的 CUDA 编程层级从三级扩展为四级：

线程（thread）→ 线程块（thread block）→ 线程块簇（thread block cluster）→ 网格（grid）

• 簇内所有 block 在物理上相邻的 SM 上并发运行，从而支持跨 SM 的高效协作。
• 使用 cluster 的最主要动机：大部分时间是为了使用分布式共享内存（DSMEM）。

1.2 为什么需要 cluster

• 每个线程块受单个 SM 上有限的共享内存与算力约束。
• 线程若想访问更多数据，就得去全局内存取——昂贵。
• 对 top-k、矩阵乘法这类需要”访问其它 SM 上的线程块数据”的算法，cluster 提供了解决方案。
• 它让算法可以用少量内存带宽换取跨多个 SM 的、显著扩大的数据可访问性。

1.3 共享内存”池化”的直觉

• 把若干线程块的共享内存汇聚起来，不同 block 的线程就能访问所需数据，而无需走全局内存。
• 但不是真的倒进一个池子：每个 block 仍然拥有自己那部分共享内存，区别只在于—— 现在线程可以访问其它 block 的共享内存。
• 这个”池”通常很小，且被限制在一个 GPC 内。

1.4 几个要点

• 线程块仍然运行在单个 SM 上——这个一对一关系没有变。
• SM 可以同时运行多个线程块——以前可以，现在依然可以。
• Cluster 是软件概念，GPC 是硬件概念。
• 没有自动分配：你必须在代码里显式定义 cluster；调度器不会自动把 block 分组。

1.5 关于 cluster 大小（重要权衡）

结论：cluster 不是越大越好；size=2 往往最优，因为更大的 cluster 带来同步开销且浪费 SM。

2. 分布式共享内存（DSMEM）

2.1 定义

DSMEM 是 H100 的特性，允许簇内不同 SM 之间直接访问共享内存。

• H100 为 cluster 实现了专用的 SM-to-SM 网络，提供对远端共享内存的快速、低延迟访问。
• 使得一个 SM 能对其它 SM 的共享内存执行 load / store / atomic 操作。
• DSMEM 可与 L2 cache 访问同时使用——应用在 SM 间通信时可叠加两条通路的总带宽。

2.2 Multicast 与 TMA + DSMEM

• 可以把 TMA 用于 DSMEM 的异步拷贝。
• DSMEM 最重要的特性之一是 multicast（多播）：把数据同时投递到多个 SM 的共享内存。
• TMA multicast 绕开了 SM-to-SM 网络的瓶颈：与其让线程显式跨 DSMEM 读写（造成拥塞与同步开销）， 不如让每个 cluster 只需一个线程发起一次 TMA multicast，把数据一次性分发给所有 SM。

3. 创建与使用线程块簇

3.1 编译期定义

在 kernel 声明里用 __cluster_dims__ 属性定义 cluster 维度：

// 编译期定义 cluster 维度（例如 2×1×1）
__global__ void __cluster_dims__(2, 1, 1) my_kernel(...) { ... }

• 定义后可以正常启动 kernel，但 grid 维度必须是 cluster 大小的整数倍。

3.2 Cluster 句柄（cooperative_groups）

#include <cooperative_groups.h>
namespace cg = cooperative_groups;

cg::cluster_group cluster = cg::this_cluster();          // 当前线程所属的 cluster 句柄
int* remote_smem = cluster.map_shared_rank(smem, target_block_rank);  // 映射到目标 block 的共享内存
remote_smem[idx] = value;                                // 跨 SM 写共享内存
unsigned int cluster_size = cluster.num_blocks();        // cluster 中 block 数
unsigned int cluster_rank = cluster.block_rank();        // 本 block 在 cluster 中的 rank

讲义提到：实际开发中不常用 cooperative_groups 这套 API，更常用 mapa（见 §8.4） 把共享内存地址转换成 cluster 内地址。

3.3 PTX 特殊寄存器（cluster 相关）

这些寄存器提供 cluster 内线程块的信息：

4. PTX 与内联 PTX

4.1 什么是 PTX

PTX 是 NVIDIA GPU 的汇编语言，是 CUDA 生态里的指令集架构（ISA）。

编译链：

CUDA C++ 源码 → 编译 → PTX → JIT 编译 → SASS → 汇编成二进制

为什么要学 PTX：它能直接访问 CUDA C/C++ 未直接暴露的底层 GPU 特性，用于手工优化性能关键段。 很多知名仓库（CUTLASS 等）大量使用 PTX，学会它就能读懂高度优化库里的各种优化技巧。

4.2 内联 PTX（Inline PTX）

内联 PTX = 用 asm 关键字把 PTX 汇编直接嵌入 CUDA C/C++。

• 好处：省去从零手写整个 PTX 文件的重活，直接在 C++ 里写 PTX 指令。
• 常配合模板：把编译期常量参数化后直接嵌入 GPU 指令。

4.3 PTX 指令的格式

asm("ptx 指令字符串" : 输出操作数 : 输入操作数 : clobbers);

• asm()：把 PTX 代码插入 CUDA 程序。
• volatile 关键字（asm volatile(...)）：防止编译器删除或移动你的 PTX 指令。
• 若不想让编译器看到你的内存访问方式，在 clobbers 里用 "memory"。

4.4 输入与输出操作数

操作数是 C++ 变量与汇编指令之间的桥梁，用约束（constraint）语法告诉编译器如何把变量映射到寄存器或内存。

• 顺序：先输出操作数 → 再输入操作数 → 最后 clobbers。
• 操作数按输出在前、输入在后统一顺序编号。
• 无输出时可留空：asm("string" :: 输入操作数)。

4.5 约束修饰符（Constraint Modifiers）

4.6 约束（Constraints）

4.7 操作数的表示

• 操作数按文本顺序用 %0, %1, %2 表示（%0 = 指令字符串后第一个变量）。
• 用约束修饰符 + 约束表示变量：输出 "+r"、"=f"、"+l"，输入 "r"、"l"、"f"。
• 花括号 {} 用于局部寄存器作用域。

4.8 指令字符串 / 输出 / 输入 / clobbers

• 指令字符串：真正要在 GPU 上执行的指令；可单行或多行（用换行符）；用 %0/%1 占位。

  asm volatile("wgmma.mma_async {%0, %1, %2, ...}");
  

• 输出操作数：在指令字符串之后；多个用逗号分隔；无输出则留空（::）。
• 输入操作数：在输出之后、以冒号分隔；只读（"n"/"r"/"f"/"l"）或读写（"+n"/"+r"/"+f"），不能只写；"+n" 可拿到编译期常量。
• Clobbers：告诉编译器”除了输出操作数之外，还有哪些资源可能被修改”， 防止编译器做错误假设而改坏代码。

5. PTX 状态空间（State Spaces）

5.1 什么是状态空间

状态空间 = GPU 线程可访问的不同内存区域，各自针对特定用途优化。

• 写 PTX 时必须指定指令作用于哪个状态空间。
• 它们不只是底层 PTX 细节——正是因为”内存放置是一等公民”，才能写出高性能代码。

5.2 各状态空间

6. 数据类型（Data Types）

6.1 基础类型

6.2 深度学习中重要的数据类型

6.3 打包数据类型（Packed Types）

把多个值打包进一个寄存器以并行操作：

这正是 H100-架构介绍.md 里”32 位是寄存器基本单位，FP16/FP8 需打包”的 PTX 层面体现。

7. 内存地址与 Shared Memory Bank

7.1 内存地址的本质

• 把 GPU 全部内存想象成一把无限长的尺子；可寻址的最小单位是字节（8 位），也叫 atom。
• 地址 = 距尺子起点的字节距离；字节地址就是尺子上的原始整数索引。
• CUDA 里很少直接操作原始字节，而是操作类型（float/int/bf16）。写 ptr + 1 时，编译器移动的是 1 个元素，而不是 1 个字节。

7.2 Shared Memory 与 Bank

为了让 warp 的 32 个线程同时访问内存，shared memory 被分成 32 个 bank，编号 0–31，每个 bank 4 字节宽。

• GPU 根据 4 字节字索引确定地址属于哪个 bank：
  • 字节 0–3 → Bank 0；字节 4–7 → Bank 1；字节 8–11 → Bank 2；字节 12–15 → Bank 3，以此类推。

7.3 Bank Conflict

• 硬件每个 bank 每 cycle 只能服务 1 个唯一地址。
  • 2 个线程冲突 → 访问串行化（2 倍慢）；最坏 32 个线程冲突 → 32 倍慢。
• 多个线程读完全相同的地址时无冲突——硬件做 multicast/broadcast，1 个 cycle 服务所有线程。
• Bank Conflict = 同一 warp 内多个线程访问”映射到同一 bank 的不同地址”。
• H100 内存控制器按 128 字节事务处理请求。
• 若 warp 请求总量超过 128 字节（如每线程加载 16 字节 float4），请求会被拆成多个事务（wave）。

7.4 Swizzling（打散）的由来

• H100 依赖 Tensor Core（读取 64×64 tile），标准线性寻址常造成大规模 bank conflict（跨步访问）。
• 例：一个 64×64 的 bf16 tile：
  • 行大小 = 64 元素 × 2 字节 = 128 字节；
  • 每行 bank 容量 = 32 bank × 4 字节 = 128 字节。
  • → 矩阵的每一列都完美地落在同一个 bank（(128/4) mod 32 = 0）。
• 我们希望 Matrix[0][0] 与 Matrix[1][0] 落在不同 bank，即使它们的步长恰好是 bank 宽度的整数倍—— 这就是 swizzle（地址打散） 的动机，后续 WGMMA 专题会用到。

8. 指针与状态空间

8.1 UVA（统一虚拟寻址）地址映射

UVA 把各状态空间映射到互不重叠的区域：

8.2 通用指针（Generic Pointers）

• 现代 CUDA（Compute Capability ≥ 3.5）里，所有指针默认是通用指针——即统一的 64 位虚拟地址空间，横跨所有状态空间。
• 没有运行时检查确定当前处于哪个状态空间。
• 可以像 CPU 指针一样做普通算术。
• 解引用时，硬件检查高位把请求路由到正确内存子系统，代价是损失几个周期做空间检测。

8.3 状态空间指针（State Space Pointers）

• 写内联 PTX（或编译器生成 PTX）时，指针可用其状态空间限定。
• 不是普通指针，只是对特定硬件指令有意义的原始位。
• 不能在 C++ 中解引用（*(float*)shared_bits 会崩溃或得到垃圾值）。
• 无运行时空间检测——硬件立刻知道它是 shared memory。
• 支持需要显式空间限定的专用指令。

8.4 cvta 与 mapa

• cvta（Convert To Address）：PTX 指令，在”通用地址”与”特定内存空间地址”之间转换指针。 CUDA C++ 对应 API：

  __cvta_generic_to_global / __cvta_generic_to_shared
  __cvta_generic_to_local  / __cvta_generic_to_constant
  

• mapa：用于把当前 CTA 的共享内存地址转换为同一 cluster 内另一个 CTA 的共享内存地址（即分布式共享内存地址转换）。
  • 关键点：mapa 接受的是 rank，不是 block ID！

8.5 Rank vs Block ID

8.6 什么时候需要这些（cvta / 状态空间指针）

• Tensor Core 的 ldmatrix 需要显式的 shared memory 地址。
• TMA 需要显式空间限定。
• Thread Block Cluster 跨 SM 共享数据需要精确的空间控制（用 cluster 专用地址）。
• MIG（多实例 GPU） 配置下，显式地址空间转换有助于保证各 GPU 实例之间的内存隔离。

9. 总结与学习衔接

9.1 四个主题速记

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

9.3 一句话记忆

Cluster 让”跨 SM 共享内存”成为可能（软件概念、限 GPC 内、size=2 常用）； 内联 PTX 让你直接指挥底层指令（asm + 约束 + clobbers）； 状态空间与指针（UVA/generic/cvta/mapa）决定了数据放在哪、怎么寻址——这是写高性能 kernel 的底层根基。

参考来源：2. Clusters, Data types, inline PTX, State Spaces.pdf（Lesson 2，47 页）。