CUDA的假象：B200单颗GPU下的NUMA代价

一颗GPU里的两颗芯片：被扁平化掩盖的NUMA拓扑

我写 GPU 代码这几年，有一个前提一直没变过：CUDA 设备是一个整体。你调用 cudaMalloc，拿到一块地址；你 launch kernel，所有 SM 都能访问这块地址。这个前提在 Hopper 时代完全成立，直到 Blackwell 把它推翻了。Vincent Zhong 在他的文章里直接说，Blackwell 把我们拽回了 NUMA 问题的战场，只不过战场从服务器机柜缩小到了一颗芯片的封装内部 来源：VZ。

B200 和 B300 物理上是两颗独立的硅 die，通过 NV-HBI 互连结合在一起 来源：VZ。每颗 die 构成一个局部性域，拥有自己的 SM 阵列和自己的 HBM 堆栈。HBM 是直接附着在自家 die 上的，访问自家 die 挂载的 HBM 比跨过去访问另一颗 die 的 HBM 明显更近。"近"这个词在延迟和功耗两个维度上都能量化。图 1 展示了每颗 die 持有一半权重行，各持一个 task counter 的布局 来源：VZ。

问题出在抽象层。CUDA 把两颗 die 呈现为一个设备、一个平坦地址空间、一个统一的 L2 缓存，你 launch kernel 时跟 Hopper 时代没有区别。但 CUDA 没有任何 API 可以查询某个 SM 属于哪颗 die 来源：VZ。这种信息隐藏让 NUMA 问题变得比传统的服务器级 NUMA 更难诊断。传统 NUMA 你至少可以用 lscpu 看看拓扑，这里你什么工具都没有，你甚至不知道自己要解决什么。开发者只能在自己运行时加探测代码，或者干脆忽略这个拓扑差异，直到 MoE 推理的性能问题暴露出来。

DeepGEMM的启动探测：用clock64给NUMA拓扑画像

DeepSeek 团队在 DeepGEMM 的 MegaMoE kernel 里率先碰上了这个问题，相关代码来自 deepseek-ai/DeepGEMM#462 来源：DG。既然 CUDA 不告诉你 SM 的归属，那就自己在启动时量出来。这个思路简单到有点粗暴，但在没有更好的 API 之前，它是最可靠的做法。

探测方法很直接：向每颗 die 放置缓冲区，让每个 SM 沿两条指针链追逐，用 clock64 计时跨 die 的延迟 来源：VZ。探针 kernel 用 __launch_bounds__(32,1) 限制每 SM 只运行一个 thread 执行探测，避免并发干扰测量。探测 kernel 用 ptx::get_sm_idx() 获取 SM 编号，用 ptx::ld_global_cg() 执行全局缓存加载指令，绕过 L1 且 L2 预先 flush，使得 cycles 计的是 HBM 往返延迟而非缓存命中 来源：VZ。每个 chunk 是 4 KB，对应 32 条 cache line，每条指向后第 7 条，确保串行依赖。

量出拓扑之后，策略是每颗 die 各持有一半权重行和独立的 task counter 来源：VZ。只有当某层的计算量达到 2 个以上 wave 时，才会触发跨 NV-HBI 的任务窃取 来源：VZ。这个门槛说明作者把跨 die 通信当作一种后备机制，只在必要时才启用。

跨die拉权重的代价：从架构设计到社区反馈

如果我忽略 NUMA 特性来写 MoE kernel，结果是从对侧 die 拉取一半的权重，同时付出延迟和功耗双重代价 来源：VZ。NV-HBI 互连的带宽虽然很快，但跨 die 访问仍有可测量的性能差距。作者用"back to solving NUMA problems"来概括核心洞察，意思是 NUMA 从未消失，只是从服务器机柜降维到了芯片内部 来源：VZ。这个洞察对正在写 Blackwell 代码的人影响很大，但对大多数人来说还是太早了。

HN 帖子截至发表时只有 1 个 point、0 条评论 来源：HN。一个关于下一代旗舰 GPU 的架构级问题，几乎无人讨论。我觉得有两个原因。一个是 DeepGEMM 的方案目前仅覆盖 MegaMoE kernel，CUDA 统一抽象层本身尚未提供显式的 NUMA 感知接口 来源：VZ，所以大多数开发者还没碰到这个问题；另一个是，NUMA 在 GPU 社区的讨论度一直不高，大家更习惯在平坦抽象层里打转。

我不确定 NVIDIA 会不会在后续 CUDA 版本中暴露 die 拓扑信息，也不确定其他推理框架是否会跟进类似的探测机制。可以确定的是，在 CUDA 抽象保持不变的前提下，Blackwell 上的性能优化如果不想丢掉跨 die 的带宽红利，就得自己写探测代码。这不是一条所有人都愿意走的路，但对于 MoE 推理这种 workload，代价可能是你实际感受得到的。对于还在用 Hopper 的开发者来说，这个趋势值得提前关注：多 die 封装架构的芯片级 NUMA 问题，未来大概率会成为通用挑战。

原文：

Blackwell Brings NUMA Back, Inside One GPU · 来源：Hacker News

作者：itfan123 值班编辑 · 发布于 2026年10月26日