LocalDmaBandwidth
本页上的每个值、偏移和地址都逐字节精确地读取自
libtpu-0.0.40-cp314wheel 中的libtpu.so（BuildID md589edbbe81c5b328a958fe628a9f2207d）。该二进制文件未被 stripped —— 下方每个符号都是反混淆后的 C++ 名称。各代带宽单元格通过解码每个 accessor 返回的 IEEE-754 立即数恢复；消费者公式则从反编译函数体加.rodata立即数解码中恢复。其他版本会不同。所有地址都是虚拟地址；对这个二进制而言，.textVMA == 文件偏移 (0xe63c000)，.rodataVMA == 文件偏移 (0x84a0000)。Itanium-ABI 注：对象的 vptr 是"vtable for X" + 0x10，因此一次虚调用*(vptr+N)会落到槽位N。
摘要
Target::LocalDmaBandwidth(MemorySpace src, MemorySpace dst) 是成本模型的片上 DMA 带宽矩阵：一张按代划分的 GB/s 数值表，覆盖 HBM、VMEM、CMEM、SMEM 之间每个 (source memory space, destination memory space) 对（以及一个 SPMEM→HBM 条目）。它不是 conv/fusion 操作数 DMA 使用的每字节周期率 —— 那条路径经过 GetBytesPerCycle 和芯片几何结构的 HBM bytes-per-second（见 memory-bandwidth-latency-model）。相反，LocalDmaBandwidth 是异步与同步 copy 比较器：唯一的消费者是三个 copy 策略决策器，它们询问“本地 DMA（例如 VMEM→VMEM rotate）是否比通过 ICI 推送相同字节更快？”并且仅当 ICI 上限优于本地 DMA 估计时才选择异步本地 copy。
熟悉 LLVM 的读者应将其映射为一个 TargetTransformInfo 成本 hook：它为 memory-space 到 memory-space 的 copy 返回一个相对吞吐量数值，由某个 transform 查询以在两种 lowering 之间决策 —— 而不是送入调度器的周期计数。该矩阵位于 Target vtable 中两个 ICI-rate accessor 之后（+0x188 ICIPerLinkDataRate，+0x190 ICIIngressEgressDataRate，随后 +0x198…+0x208 是 LocalDmaBandwidth 单元格，+0x210 SpmemToHbm），而 LocalDmaBandwidth 本身是一个纯 (src,dst)→slot 路由器，会尾调用匹配的 accessor，或在没有匹配时返回空 optional。
本页端到端记录该矩阵：dispatch 路由器和 (src,dst)→slot 映射、从每个 accessor 转录的完整按代 GB/s 数值（Dragonfish、Pufferfish、Viperfish std/lite、Ghostlite —— 基类 Target 返回 0）、异步 copy 消费者（UseAsyncDataCopy、ShouldUseAsyncLocalCopy、SparseCore AllGather 策略）及其 min(ICIIngressEgress, 2·ICIPerLink) ICI 上限，并且 —— 作为对照 —— 记录独立的每字节周期定价器（GetBytesPerCycle/WindowCycles/DefaultHbmInitLatency），即操作数 DMA 调度入口实际消费的部分，以及它按代的 InitialDmaLatencyInNs 启动常量。
对于重新实现，契约是：
- dispatch 路由器
LocalDmaBandwidth(src,dst)—— 它的(src,dst)→vtable-slot映射、MemorySpace枚举{HBM=1, VMEM=3, CMEM=4, SMEM=5}，以及未命中时的空 optional 返回。 - 完整的按代 GB/s 矩阵（转录如下），以及 Viperfish
variant_name == "lite"(0x6574696c) 的 std/lite 分支。 - 异步 copy 消费者：
local_cost = LocalDmaBandwidth(src,dst) · (n−1)对比ici_ceiling = min(ICIIngressEgress, 2·ICIPerLink)；当且仅当ici_ceiling ≤ local_cost时异步。 - 与每字节 MemXfer 定价器的对照：
bytes_per_cycle = HbmFullChipBytesPerSecond / (TC_freq_MHz·1e6) / CoresPerChip，WindowCycles = transfer_bytes / bytes_per_cycle + init，DefaultHbmInitLatency = InitialDmaLatencyInNs · (TC_freq_MHz/1000)cycles —— 以及按代的InitialDmaLatencyInNs（240 / {555,50} / 1200,0 ns）。
| Dispatch 路由器 | Target::LocalDmaBandwidth(MemorySpace,MemorySpace) @0x1d6168e0 |
| MemorySpace 枚举 | HBM=1, VMEM=3, CMEM=4, SMEM=5（SPMEM 通过自己的槽位） |
| 矩阵 vtable 范围 | +0x198 (HbmToHbm) … +0x208 (SmemToSmem); +0x210 SpmemToHbm 尾部 |
| 基类默认值 | Target::LocalDmaBandwidth* accessors @0x1d48fa00… 全部返回 0 |
| 消费者（真实） | (anon)::UseAsyncDataCopy @0x1380a480; ShouldUseAsyncLocalCopy @0x133eff40 |
| ICI 上限 | min(ICIIngressEgress[+0x190], 2·ICIPerLink[+0x188]) |
| 每字节定价器（独立） | fusion_util::GetBytesPerCycle @0x1454dd00 → WindowCycles @0x14552660 |
| 调度入口 | cost_model_util::RecordMemXferCyclesImpl @0x13844e80 (R9/R10, R11/R12) |
Dispatch 路由器 — LocalDmaBandwidth
目的
LocalDmaBandwidth(src, dst) @0x1d6168e0 是一个纯路由器：给定两个 MemorySpace 枚举值，它计算一个 vtable 字节槽位并尾调用对应的成对 accessor（LocalDmaBandwidthHbmToHbm，……）。对于无法识别的 pair，它返回空 optional {value=0, has_value=0}。路由器本身不保存任何值；GB/s 数值位于它 dispatch 进入的按代 accessor 叶子中。
算法
c
function LocalDmaBandwidth(Target* this, uint8 src, uint8 dst): // @0x1d6168e0
// The decompiler renders the slot as a guarded cascade of XOR tests
// (src^1 == HBM, src^3 == VMEM, src^4 == CMEM, src^5 == SMEM); the first
// matching (src,dst) pair leaves `slot` at the table value below.
slot = (src,dst) -> {
HBM ->HBM : 0x198, HBM ->VMEM : 0x1a0, HBM ->SMEM : 0x1a8,
VMEM->HBM : 0x1b0, VMEM->VMEM : 0x1b8, VMEM->CMEM : 0x1c0, VMEM->SMEM : 0x1c8,
CMEM->HBM : 0x1d0, CMEM->VMEM : 0x1d8, CMEM->CMEM : 0x1e0, CMEM->SMEM : 0x1e8,
SMEM->HBM : 0x1f0, SMEM->VMEM : 0x1f8, SMEM->CMEM : 0x200, SMEM->SMEM : 0x208,
else : MISS
}
if slot == MISS: return optional{ value=0, has_value=false }
return (*(vtable + slot))(this) // tail-call the accessor
```text
`MemorySpace` 枚举从 XOR 比较数恢复：`HBM=1`、`VMEM=3`、`CMEM=4`、`SMEM=5`。SPMEM 不在这个路由器中；单独的 `LocalDmaBandwidthSpmemToHbm` accessor 位于 vtable `+0x210`（紧跟 `+0x208` 的 `SMEM->SMEM` 单元格之后），并且会被直接到达（SparseCore 路径），而不是通过这个 `(src,dst)` dispatch。
> **注意 —** 两个 ICI-rate accessor 是矩阵在 `Target` vtable 中的相邻项：`+0x188 = ICIPerLinkDataRate`、`+0x190 = ICIIngressEgressDataRate`，随后是从 `+0x198`（`HBM->HBM`）到 `+0x208`（`SMEM->SMEM`）的 LocalDmaBandwidth 单元格，`SpmemToHbm` 位于 `+0x210`。这种相邻并非装饰 —— 同一个消费者（`UseAsyncDataCopy`）读取 LocalDmaBandwidth 单元格和两个 ICI 速率来做一次决策，因此编译器将它们连续布局。
### 函数映射
| 函数 | 地址 | 作用 |
|---|---|---|
| `Target::LocalDmaBandwidth(MemorySpace,MemorySpace)` | `0x1d6168e0` | `(src,dst)→slot` 路由器；未命中 → `{0,false}` |
| `Target::LocalDmaBandwidthHbmToHbm`（基类） | `0x1d48fa00` | 基类默认值 —— 返回 `0`（每代会覆盖） |
| `Target::ICIPerLinkDataRate`（vtable `+0x188`） | 按代 | 每条 SerDes link 的 ICI 速率（消费者上限项） |
| `Target::ICIIngressEgressDataRate`（vtable `+0x190`） | 按代 | 完整的双向芯片 ICI 聚合 |
---
## 按代带宽矩阵
每个 accessor 返回一个 IEEE-754 `double`（GB/s）。基类 `Target` accessors（`@0x1d48fa00…`）全部 `return 0` —— 在派生的按代 `Target` 覆盖 vtable 槽位之前，该矩阵没有意义。下表转录了从每个按代 accessor 解码出的立即数。每个单元格都通过解码 accessor 返回的 bit-pattern 确认。
### Ghostlite (v6e) — `@0x1d4973c0 …`
| (src→dst) | Accessor `@addr` | bit-pattern | GB/s |
|---|---|---|---|
| HBM→HBM | `0x1d4973c0` | `0x4050000000000000` | 64 |
| HBM→VMEM | `0x1d4973e0` | `0x4094140000000000` | 1285 |
| HBM→SMEM | `0x1d497400` | `0x404B800000000000` | 55 |
| VMEM→HBM | `0x1d497420` | `0x4096600000000000` | 1432 |
| VMEM→VMEM | `0x1d497440` | `0x4050000000000000` | 64 |
| VMEM→SMEM | `0x1d497460` | `0x404B800000000000` | 55 |
| SMEM→HBM | `0x1d497480` | `0x404B800000000000` | 55 |
| SMEM→VMEM | `0x1d4974a0` | `0x404B800000000000` | 55 |
| SMEM→SMEM | `0x1d4974c0` | `0x403C000000000000` | 28 |
| SPMEM→HBM | `0x1d4974e0` | `0x4082600000000000` | 588 |
### Viperfish (v5p std / v5e lite) — `@0x1d49a320 …`
每个 Viperfish accessor 都会基于 variant 字符串分支：它读取 `this+951`（`+0x3b7`）处的 libc++ `std::string variant_name` SSO 标志字节、其长度（SSO inline 或 `this+0x3a8` heap length），并将 `this+0x3a0` 处的 4 字节 payload 与 `0x6574696c`（`"lite"`，little-endian）比较。4 字符 `"lite"` 匹配时返回 **lite (v5e)** 值；任何其他字符串都会落到 **std (v5p)** 值。
| (src→dst) | Accessor `@addr` | std (v5p) | lite (v5e) |
|---|---|---:|---:|
| HBM→HBM | `0x1d49a320` | 72 | 308 |
| HBM→VMEM | `0x1d49a380` | 1198 | 822 |
| HBM→SMEM | `0x1d49a3e0` | 55 | 56 |
| VMEM→HBM | `0x1d49a440` | 1224 | 828 |
| VMEM→VMEM | `0x1d49a4a0` | 72 | 827 |
| VMEM→SMEM | `0x1d49a500` | 55 | 56 |
| SMEM→HBM | `0x1d49a560` | 55 | 56 |
| SMEM→VMEM | `0x1d49a5c0` | 55 | 56 |
| SMEM→SMEM | `0x1d49a620` | 28 (both) | 28 (both) |
| SPMEM→HBM | `0x1d49a640` | 587.4 (both) | 587.4 (both) |
> **易错点 —** variant 分支是在每个 accessor 中*内联计算*的，而不是通过数值 variant index。单 TC 的 v5e die 一致报告*更高*的片内 loopback 带宽（VMEM→VMEM 827，而 std die 为 72），因为 lite 部件没有第二个 core 与本地 DMA fabric 竞争；反过来，它报告*更低*的 HBM-bound 带宽（HBM→VMEM 822，而 1198），反映其减半的 HBM stack。如果重新实现时把一个 variant 的单元格用于两者，就会错误定价另一个 variant 上的异步 copy 决策。
### Pufferfish (v4) — `@0x1d494340 …`
Pufferfish 是唯一拥有一等 CMEM 的代，因此它的矩阵最宽（15 个单元格）。
| (src→dst) | Accessor `@addr` | bit-pattern | GB/s |
|---|---|---|---:|
| HBM→HBM | `0x1d494340` | `0x407E000000000000` | 480 |
| HBM→VMEM | `0x1d494360` | `0x407E100000000000` | 481 |
| HBM→SMEM | `0x1d494380` | `0x4041000000000000` | 34 |
| VMEM→HBM | `0x1d4943a0` | `0x40915C0000000000` | 1111 |
| VMEM→VMEM | `0x1d4943c0` | `0x4081000000000000` | 544 |
| VMEM→CMEM | `0x1d4943e0` | `0x4091840000000000` | 1121 |
| VMEM→SMEM | `0x1d494400` | `0x4041000000000000` | 34 |
| CMEM→HBM | `0x1d494420` | `0x4090E00000000000` | 1080 |
| CMEM→VMEM | `0x1d494440` | `0x40A2460000000000` | 2339 |
| CMEM→CMEM | `0x1d494460` | `0x4092A40000000000` | 1193 |
| CMEM→SMEM | `0x1d494480` | `0x4041000000000000` | 34 |
| SMEM→HBM | `0x1d4944a0` | `0x4041000000000000` | 34 |
| SMEM→VMEM | `0x1d4944c0` | `0x4041000000000000` | 34 |
| SMEM→CMEM | `0x1d4944e0` | `0x4041000000000000` | 34 |
| SMEM→SMEM | `0x1d494500` | `0x4031000000000000` | 17 |
### Dragonfish (v3) 和基类 Target
Dragonfish 只覆盖两个单元格；其他每个槽位都会落回基类 `Target` accessor（返回 `0`）。Jellyfish (v2) 在此 build 中完全不提供任何 `LocalDmaBandwidth*` 覆盖 —— 它的异步 copy 路径从不查询该矩阵。
| Gen | (src→dst) | Accessor `@addr` | GB/s |
|---|---|---|---:|
| Dragonfish (v3) | HBM→VMEM | `0x1d48fa20` | 423 |
| Dragonfish (v3) | VMEM→HBM | `0x1d48fa60` | 423 |
| base `Target` | all cells | `0x1d48fa00…0x1d48fbe0` | 0 |
> **怪癖 —** SMEM 单元格聚集在一小组重复常量上（Pufferfish 对每个涉及 SMEM 的 pair 都是 `34`，除了 SMEM→SMEM = `17`；Ghostlite 对每个 SMEM pair 都是 `55`，除了 SMEM→SMEM = `28`；Viperfish std `55` / lite `56`）。SMEM（标量寄存器窗口）是一块很小、很慢的内存；成本模型将每个涉及它的传输都视为单个低固定速率，而不是对 source/destination pairing 建模。只有 HBM/VMEM/CMEM 单元格携带按 pair 区分的数值。
---
## 真实消费者 — 异步与 ICI Copy 决策
### 目的
该矩阵恰好有三个调用者，全部是 copy 策略决策器：`(anon)::UseAsyncDataCopy` `@0x1380a480`、`ShouldUseAsyncLocalCopy` `@0x133eff40`，以及 SparseCore AllGather 策略 `TcStyleSinglePhaseAGTransferStrategy::ComputeUseAsyncLocalCopy` `@0x1335f2e0`。它们没有任何一个把周期计数送入调度器；每个都将本地 DMA 带宽与 ICI 上限比较并返回一个布尔值 —— “使用异步本地 copy”还是“走 ICI”。
### 算法
```c
function UseAsyncDataCopy(n_elems, …, Target& t, fn src_of, fn dst_of): // @0x1380a480
if early_flags: return 0
if n_elems <= 0: return 1 // trivial copy is always async-OK
src = src_of(0); dst = dst_of(0)
if not t.SupportsDmaMode(src, dst, …)[vtable+0x2a0]: return 0 // gen must allow the DMA
if not (t.ICIPerLinkDataRate[+0x188] && t.ICIIngressEgressDataRate[+0x190]):
return 0
local_bw = t.LocalDmaBandwidth(src, dst) // GB/s for this (src,dst) pair
local_cost = local_bw · n_elems // (var_38 = n_elems as double)
ici_ceiling = min( t.ICIIngressEgressDataRate, // r13
2 · t.ICIPerLinkDataRate ) // 2·rbx
// async iff the ICI ceiling does NOT beat the local-DMA estimate
return (ici_ceiling <= local_cost) // vucomisd @0x1380a5ad比较是一条 vucomisd ici_ceiling, local_cost：当 ICI 上限小于或等于本地 DMA 带宽（按元素数缩放）时，也就是本地 DMA 有竞争力时，函数选择异步本地 copy。SupportsDmaMode（vtable +0x2a0）为整个决策设门槛 —— 不支持请求的 (src,dst) DMA mode 的代会立即回退到同步/ICI 路径。
注意 —
ICIPerLinkDataRate是一条 SerDes link 的带宽；ICIIngressEgressDataRate是完整的双向芯片聚合。上限min(IngressEgress, 2·PerLink)建模两个 regime：通过单对 link 的传输受限于2·PerLink（一进一出），而能饱和所有 link 的传输受限于全芯片 IngressEgress 聚合。因此该决策是“本地片上 DMA 是否快于单 link pair 和 all-links ICI 中更好的那个？”
两个带宽来源 — 不要混淆
这是关于 LocalDmaBandwidth 最重要的结构性事实：它不是操作数 DMA 调度器定价所用的带宽。成本模型有两个不同的片上带宽来源，它们有不同消费者和不同单位：
| 来源 | 存储位置 | 消费者 | 单位 | 用途 |
|---|---|---|---|---|
LocalDmaBandwidth(src,dst) 矩阵 | 按代 vtable +0x198… | UseAsyncDataCopy / ShouldUseAsyncLocalCopy / SparseCore AG | GB/s（相对） | 异步与 ICI copy 策略决策 |
HbmFullChipBytesPerSecond (chip_parts) | Target+0x4f0 | GetBytesPerCycle → WindowCycles | bytes/sec → bytes/cycle | 每字节操作数 DMA 周期沉积（R10/R12） |
如果重新实现把操作数 DMA 周期定价路由到 LocalDmaBandwidth 矩阵，就会是错的：该矩阵是比较器 heuristic；每字节成本使用芯片几何结构的 HBM bytes-per-second。下一节记录第二个来源，作为对照并锚定调度入口。
每字节定价器（独立路径）— GetBytesPerCycle / WindowCycles
目的
操作数 DMA 的周期成本 —— 最终落入调度器 ResourceVector 的值 —— 由 MemXfer 定价器计算，后者从不触碰 LocalDmaBandwidth 矩阵。GetBytesPerCycle 从芯片几何结构推导每 core、每 cycle 的字节预算；WindowCycles 用传输字节数除以它；DefaultHbmInitLatency 加上一次性支付的 DMA 启动。字节数本身来自 WindowDescription 字节模型。
算法
c
function GetBytesPerCycle(HloInstruction* inst, Target& t, MemorySpace ms): // @0x1454dd00
CHECK(ms == HBM || ms == CMEM) // else Fatal (fusion_util.cc:2354)
freq_hz = t.TensorCoreFrequencyInMegaHertz() · 1e6 // const @0xa2e0208 = 1,000,000.0
bw_Bps = (ms == HBM) ? t.HbmFullChipBytesPerSecond() // Target+0x4f0 (= *(this+158))
: t.CmemFullChipBytesPerSecond() // Target+0x4f8 (= *(this+159))
bpc = bw_Bps / freq_hz / t.CoresPerChip(coreType) // BYTES PER TENSORCORE CYCLE (per core)
// HBM-only override: a positive TpuCompEnv[+0x1040] or BackendConfig field-14
// value REPLACES bpc (flag-driven; default uses the computed value).
return bpc
function WindowCycles(MemUnit, wd, Target& t, bytes_per_cycle, …, c): // @0x14552660
if MemUnit == t.MemUnitFromKiB(0): return 0.0 // sentinel — no transfer
init = (c != -1) ? (t.TensorCoreFrequencyInMegaHertz()/1000.0)·c
: DefaultHbmInitLatency(MemUnit, wd, t)
count_desc = (t.vtable[+0x590]() == 1) // per-gen count-descriptors predicate
base = WindowCyclesGenericTargetAgnostic(MemUnit, wd, count_desc, bytes_per_cycle) + init
if t[+0x398] >= 5: // v5p+ (Target+0x398 = TpuVersion)
// two-direction {HBM, VMEM} blend: divide the {dir0,dir1} byte pair by 1000.0
// (xmmword @0xa2ce650 = [1000.0, 1000.0]), multiply by TC freq, take per-dir MAX,
// gated by a percentage > 0.01 test (xmmword @0xa2d8020 = [0.01, 0.01])
return max(dir0_cycles, dir1_cycles) + init // @0x1455282e
else:
return max(dir0_cycles, dir1_cycles) + init // simple add path (@0x14552888)
function DefaultHbmInitLatency(MemUnit, wd, Target& t): // @0x14552ca0
elem = (BYTE[*(wd+0x20) + 0xb] >> 2) & 0x1f // operand-shape element type
init_ns = t.vtable[+0x20].InitialDmaLatencyInNs(MemUnit, window_bytes)
return init_ns · (t.TensorCoreFrequencyInMegaHertz() / 1000.0) // ns · GHz = cycles
// const @0xa2e0430 = 1000.0
```text
`GetBytesPerCycle` 是单位锚点：*周期定价器中*的 `LocalDma`/HBM 带宽是**每 TensorCore cycle、每 core 的 bytes**，按 `(bytes/sec) ÷ (TC_freq_MHz·1e6 cycles/sec) ÷ cores_per_chip` 推导。时钟是 TensorCore 时钟（`Target+0x90c`），不是单独的 DMA 时钟。`HbmFullChipBytesPerSecond` 是 `*(this+158) == this+0x4f0`；`CmemFullChipBytesPerSecond` 是 `*(this+159) == this+0x4f8` —— 两者都由 `Target::Init` 中的 chip_parts `bytes_per_second × stack_count` 折叠而来。
`WindowCycles` 的 v5p+ 分支（`Target+0x398 >= 5`）混合两个传输方向：它加载 `{HBM, VMEM}` fragment-byte pair，将每个除以 `1000.0`，乘以 TC frequency，通过 `> 0.01` percentage 测试为每个方向设门槛，并取逐方向最大值。`fusion_util.cc:3474` 处的 VLOG-6 trace 明确命名了各 lane —— `hbm_percentage`、`vmem_percentage`、`hbm_bytes`、`vmem_bytes`、`hbmbw`、`vmembw`、`hbmlatency`、`vmemlatency`、`next_gen` —— 确认 v5p+ 模型是双 lane（HBM vs VMEM）max，而不是单次传输。
### 按代 `InitialDmaLatencyInNs`（vtable `+0x20`）
一次性支付的 DMA 启动延迟，以纳秒为单位，`DefaultHbmInitLatency` 会将其缩放为 cycles。每个按代 accessor 都经过逐字节精确解码；`.rodata` 立即数直接从 `.so` 解码。
| Gen | Accessor `@addr` | rule | non-special | special |
|---|---|---|---:|---:|
| v2 Jellyfish | `0x1d48f3a0` | constant for all spaces (`@0xa2de6a8`) | 240 ns | — |
| v4 Pufferfish | `0x1d493d00` | `table[ms==CMEM]` (`@0xa2dcd40`) | 555 ns (HBM) | 50 ns (CMEM) |
| v5p/v5e Viperfish | `0x1d499ca0` | `(ms==VMEM) ? 0 : 1200` (`@0xa2df9f8`) | 1200 ns | 0 (VMEM) |
| v6e Ghostlite | `0x1d496dc0` | `(ms==VMEM) ? 0 : 1200` (`@0xa2df9f8`) | 1200 ns | 0 (VMEM) |