【infra之路】Shared Memory 的访问范围：严格限定于单个 SM

原创于 2026-08-22 22:27:56 发布 · 123 阅读

5 ·

本内容遵循CC 4.0 BY-SA版权协议

GEO检测

收录于

infra

Shared Memory 是 SM 私有的，只有分配到该 SM 上的 Block 才能访问。 其他 SM 上的线程无法直接读写任何非本 SM 的 Shared Memory。

核心规则

问题	答案
同一 SM 内的多个 Block 能否共享 SMEM？	❌ 不能。每个 Block 有独立的 SMEM 分配，互不可见
不同 SM 上的线程能否访问同一块 SMEM？	❌ 绝对不能。SMEM 没有全局地址空间
SMEM 的生命周期	从 Block 调度到 SM 开始，到 Block 执行结束为止
SMEM 的物理位置	嵌入在 SM 内部，与 L1 Cache 共用 SRAM，片外不可达

为什么这样设计？

┌─────────────────────────────────────────────────────┐
│                    GPU Die                          │
│                                                     │
│  ┌──────────┐  ┌──────────┐       ┌──────────┐    │
│  │   SM 0   │  │   SM 1   │ ...   │   SM N   │    │
│  │ ┌──────┐ │  │ ┌──────┐ │       │ ┌──────┐ │    │
│  │ │ SMEM │ │  │ │ SMEM │ │       │ │ SMEM │ │    │
│  │ │(私有)│ │  │ │(私有)│ │       │ │(私有)│ │    │
│  │ └──────┘ │  │ └──────┘ │       │ └──────┘ │    │
│  └──────────┘  └──────────┘       └──────────┘    │
│        ↕              ↕                  ↕         │
│  ┌─────────────────────────────────────────────┐   │
│  │           L2 Cache (全芯片共享)              │   │
│  └─────────────────────────────────────────────┘   │
│                       ↕                            │
│  ┌─────────────────────────────────────────────┐   │
│  │              HBM / GDDR (全局)               │   │
│  └─────────────────────────────────────────────┘   │
└─────────────────────────────────────────────────────┘

带宽：SMEM 提供 ~20TB/s 级带宽，靠的是每个 SM 独立一份。如果全局共享，带宽会坍缩到 HBM 级别（~3TB/s）
延迟：SMEM 延迟 ~30 cycles，因为物理距离在 SM 内部。跨 SM 访问不可能保持这个延迟
同步：__syncthreads() 只同步同一 Block 内的线程，正是因为 SMEM 的作用域就是 Block

如果需要跨 SM 通信怎么办？

必须通过更高层级的存储中转：

需求	机制	延迟代价
同 Grid 内 Block 间通信	Global Memory + `__threadfence_system()` / Cooperative Groups	~400+ cycles
同 Kernel 内流水线阶段	Persistent Thread + Global Memory 信号量	~400+ cycles
跨 Kernel 数据传递	Global Memory / Managed Memory	Kernel launch overhead
同 SM 内 Warp 间通信	Shared Memory ✅（同一 Block 内）	~30 cycles

⚠️ 注意：即使是同一个 Grid 中的两个 Block 被调度到同一个 SM 上，它们的 SMEM 仍然是隔离的。SM 会为每个活跃 Block 分配独立的 SMEM 区域，彼此不可见。

常见误解澄清

误解	事实
“同一 SM 上的 Block 可以共享 SMEM”	❌ 每个 Block 有独立 SMEM 分区，硬件强制隔离
“SMEM 可以通过指针传给其他 SM”	❌ SMEM 地址是 SM 局部虚拟地址，跨 SM 无意义
“`__shared__` 变量在所有 SM 上是同一份”	❌ 每个 SM 上每个 Block 都有自己的一份副本
“Hopper 的 TMA 让 SMEM 变全局了”	❌ TMA 只是异步拷贝引擎，目标仍是本 SM 的 SMEM