H200// Clusters
← H200 速查主页 / 特性卡片墙

H200 · HOPPER GH100 · 特性专题

Thread Block Clusters

★ Hopper 集群主线 · CUDA 编程层级新加一层，是 DSMEM / TMA multicast 的前提

01

是什么

DEFINITION

thread block cluster 是 Hopper 在 thread block 与 grid 之间新增的一级——一组保证同驻于同一 GPC（内含多个 SM 的硬件分区）的协作 CTA，可彼此同步、可直接访问彼此的 shared memory。

NOTE · 编程层级（Hopper 新增 cluster 层）

thread → warp（32）→ warp group（128，wgmma 单位）→ thread block（单 SM）→ thread block cluster（≤8 SM，同 GPC）★ → grid

02

为什么需要（vs A100）

RATIONALE · VS A100

A100 的 thread block 相互独立，block 间协作只能靠 global memory + grid-level sync（this_grid().sync()），数据要经 L2 往返、粗粒度。cluster 提供细粒度、低延迟的跨 SM 协作：

03

关键数字

KEY NUMBERS · CONSTRAINTS
项值注
Hopper 最大 cluster8 个 thread blockcluster_dims.x·y·z ≤ 8（16 是 Blackwell）
cluster 必须同驻于同一 GPC物理邻近，专用 SM-to-SM 网络
grid 与 cluster 关系grid 各维 = cluster 各维整数倍否则 launch 失败
声明方式编译期 __cluster_dims__(X,Y,Z)或运行期 cudaLaunchKernelEx + cudaLaunchAttributeClusterDimension
DSMEM 可寻址范围同 cluster 全部 block 的 smem即 ≤8 SM 的 smem
cluster 同步指令barrier.cluster.arrive / .waitsm_90+ 特殊寄存器 %cluster_ctarank
DSMEM vs 走 global~快 7×NVIDIA 官方口径
04

可编译示例

COMPILABLE EXAMPLE · VERBATIM

逐字摘自 CUTLASS cluster_sm90.hpp + PTX §9.7.14.3。

// 编译期声明 cluster 维度（2 blocks 沿 X）；grid 各维须是其整数倍
__cluster_dims__(2, 1, 1)
__global__ void cluster_kernel(int* out) {
  uint32_t rank;
  asm volatile("mov.u32 %0, %%cluster_ctarank;\n" : "=r"(rank) :);  // cluster 内 block 序号

  // cluster.sync() == arrive; wait; （硬件 barrier，跨 SM）
  asm volatile("barrier.cluster.arrive.aligned;\n" : :);
  asm volatile("barrier.cluster.wait.aligned;\n"   : :);

  out[rank] = (int)rank;
}

/* 运行期声明（与编译期二选一，但数值必须一致）：
   cudaLaunchAttribute attr{};
   attr.id = cudaLaunchAttributeClusterDimension;
   attr.val.clusterDim = {2,1,1};
   cudaLaunchConfig_t cfg{...};
   cfg.attrs = &attr;  cfg.numAttrs = 1;
   cudaLaunchKernelEx(&cfg, cluster_kernel, out);        */
05

注意事项 / 陷阱

PITFALLS · CHECKLIST
CAUTION · 陷阱清单
  • __cluster_dims__ 乘积必须 ≤ 8（运行期可查询 cudaOccupancyMaxActiveClusters 获实际可同驻上限，常更低）。
  • 编译期 __cluster_dims__ 与运行期 attribute 必须一致，不匹配 → launch 失败。
  • Clusters 是 DSMEM 与 TMA multicast 的前提；不开 cluster 却用这两个特性是 UB。
  • 同驻要求 → 高 smem / 寄存器需求会压低 cluster occupancy，过大 cluster + 重资源可能直接 launch 失败。
  • barrier.cluster.arrive 只发到达信号不阻塞，必须配 .wait 才真正同步；DSMEM 读前少任一个即数据竞争。
06

数据源

SOURCES