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

H200 · HOPPER GH100 · 特性专题

DSMEM · 分布式共享内存

★ Hopper 共享主线 · 集群内 SM 间直访彼此 smem，不经 L2/global

01

是什么

DEFINITION

DSMEM 让 cluster 内任一线程通过统一分布式地址空间直接寻址同 cluster 其它 CTA 的 shared memory，延迟接近本地 smem，无需 L2 / global 往返。同 cluster 所有 CTA 的 smem 被映射进一个逻辑地址空间。

02

为什么需要（vs A100）

RATIONALE · VS A100

无 cluster 时，block 0 读不到 block 1 的 smem——要交换数据必须 write 到 global → sync → 再从 global read，经 L2 两个来回。DSMEM 把这条路径压缩成一次直连 load/store。典型用途：

03

关键数字

KEY NUMBERS · LATENCY
项值注
可寻址范围同 cluster 全部 block 的 smem≤ 8 SM
sm_90a 每 SM smem 上限228 KB静态声明上限 48KB，扩展需 opt-in
DSMEM 访问延迟（实测 H800）~200–210 cyc本地 smem ~33cyc 的 ~6×；global ~650cyc 的 ~1/3
vs 走 global 交换~快 7×NVIDIA 官方口径
地址映射指令mapa.shared::cluster把本地 smem 地址映射到指定 rank 的 CTA
远端 load/store 状态空间.shared::cluster默认 .shared == .shared::cta
04

可编译示例

COMPILABLE EXAMPLE · VERBATIM

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

// 把本地 smem 地址 a 映射成 cluster 内 CTA `rank` 的同偏移地址
__device__ __forceinline__ uint32_t set_block_rank(uint32_t smemAddr, uint32_t rank) {
  uint32_t result;
  asm volatile("mapa.shared::cluster.u32 %0, %1, %2;\n"
               : "=r"(result) : "r"(smemAddr), "r"(rank));
  return result;
}

__cluster_dims__(2,1,1) __global__ void dsmem_kernel(int* out) {
  __shared__ int tile;
  uint32_t local = __cvta_generic_to_shared(&tile);
  tile = (int)blockIdx.x;
  __syncthreads();                                       // 本 block 写对内可见

  asm volatile("barrier.cluster.arrive.aligned;\n" : :); // cluster 同步
  asm volatile("barrier.cluster.wait.aligned;\n"   : :);

  uint32_t neighbor_rank = 1 - blockIdx.x;               // 2-block cluster: 0<->1
  uint32_t remote = set_block_rank(local, neighbor_rank); // mapa 到邻居
  int v;
  asm volatile("ld.shared::cluster.u32 %0, [%1];\n" : "=r"(v) : "r"(remote));
  if (threadIdx.x == 0) out[blockIdx.x] = v;          // 读到邻居的值
}
05

注意事项 / 陷阱

PITFALLS · CHECKLIST
CAUTION · 陷阱清单
  • DSMEM 必须在 cluster launch 下；隐式 1×1×1 cluster 无 peer，远端寻址 UB。
  • 远端 smem 读只在双重同步后可见：__syncthreads()（本地）加 barrier.cluster（集群）——少任一即数据竞争。
  • 每 block 的 smem 大小计入同驻预算，大 smem 会压低 cluster 上限 / occupancy。
  • mapa 的 rank 是 1-D %cluster_ctarank，不是 3-D %cluster_ctaid；2D/3D cluster 要自己算线性 rank。
  • 远端 st.async / TMA-into-neighbor 必须给远端 mbarrier 发信号（barrier 地址也要 mapa），本地 barrier 不被通知 → 死锁。
06

数据源

SOURCES