Internal Wiki

主题

内存带宽与延迟模型 ​

本页中的每个偏移、值和地址都从 libtpu-0.0.40-cp314 wheel 中的 libtpu.so 按字节精确读取（BuildID md5 89edbbe81c5b328a958fe628a9f2207d）。该二进制文件未 strip，下面的每个符号都是 demangle 后的 C++ 名称。其他版本会不同。所有地址都是虚拟地址；对这个二进制文件，.rodata VMA == 文件偏移（0x84a0000），这里引用的 rodata double 是从这些 VA 处的 .rodata 解码而来。

摘要 ​

本页记录内存侧成本模型：TPU 成本模型中用于按通过某个内存层级移动 N 字节需要多少 TensorCore 周期来给 DMA 定价的那一半。WindowDescription 字节成本把窗口化访问转换为字节数，而本页把该字节数转换为周期。该模型恰好有两项，并求和：

text
dma_cycles = init_latency_cycles  +  transferred_bytes / bytes_per_cycle
             └─ fixed startup ─┘     └────── bandwidth-proportional ──────┘

启动项是一个按（目标层级）计的固定常量，即 InitialDmaLatencyInNs 虚函数，并由 DefaultHbmInitLatency 缩放到周期。带宽项是 bytes / bytes_per_cycle，其中 bytes_per_cycle 由整芯片 HBM（或 CMEM）字节速率除以 TensorCore 时钟和核心数一次性推导得到。这两项被放入两个独立的 ResourceVector lane，调度器对这些 lane 取最大值：当固定启动占主导时，传输是延迟受限；当 bytes / bytes_per_cycle 占主导时，传输是带宽受限。这种逐 lane 的 max 就是该模型的 roofline。在 v5p 及更新版本上，WindowCycles 自身内部还有第二个独立的双 lane 计算：带宽会被定价两次，一次针对 HBM 字节数，一次针对 VMEM 字节数，两者中较慢者胜出；但该块是可选启用的（只有卷积 classic-window 路径同时提供这两个字节数），通用的已定价 DMA 会让它保持休眠。

熟悉的参照框架是教科书式 roofline：time = max(latency, size / bandwidth)。重新实现者必须遵守的差异是：(1) 单位是 TensorCore 周期，不是秒，这两项都预先乘以 TC 频率，因此能与 CycleTable 吞吐整数复合；(2) 启动延迟按每个 resource lane收费一次，而不是按每次传输收费，因此进入同一个 lane 的一串 DMA 只支付一次启动成本；(3) 按层级的速率常量是按代际虚函数覆盖返回的 std::optional<double>，缺失值表示“这个层级对未建模”，而不是“零带宽”。

对于重新实现，契约是：

  • 按代际的每（dst 层级）InitialDmaLatencyInNs 常量（HBM/VMEM/SMEM/CMEM），以及它们三种不同的分派形态（常量 / 2 项表 / VMEM-vs-other 分支 / LogFatal 基类）。
  • DefaultHbmInitLatency = InitialDmaLatencyInNs(elem, bytes) · (TC_MHz / 1000.0)，即 ns→周期转换，以及短路无传输 MemUnit 的 MemUnitFromKiB(0) 哨兵。
  • 字节→周期转换：MemUnitFromKiB（KiB→granule-unit 移位）和 ChunkGranules（(ChunkCellCount·4) / ChunkGranuleBytes），即 windowing_util::Size 向上取整到的 granule。
  • bytes_per_cycle = FullChipBytesPerSecond / (TC_MHz · 1e6) / CoresPerChip，来自 GetBytesPerCycle，包括仅 kHbm/kCmem 的 CHECK 和两个 env/backend-config 覆盖。
  • WindowCycles 中的 roofline 组装：启动 init 加法、围绕按方向 HBM/VMEM max 的 (hbm_b | vmem_b) 可选门控，以及内部 v5p+（Target+0x398 >= 5）门控。
  • 供给复制模式决策（ShouldUseAsyncLocalCopy）的按层级 LocalDmaBandwidth<Src>To<Dst> GB/s 矩阵，它不同于周期成本 bytes_per_cycle。
启动延迟（周期）fusion_util::DefaultHbmInitLatency(MemUnit, WindowDescription, Target) @0x14552ca0
按层级启动（ns）Target::InitialDmaLatencyInNs(MemorySpace, double) @0x1d61c880（基类 LogFatal）
Roofline 组装fusion_util::WindowCycles(MemUnit, WindowDescription, Target, double, long, long, long) @0x14552660
字节/周期预算fusion_util::GetBytesPerCycle(HloInstruction, Target, MemorySpace) @0x1454dd00
KiB→granule 转换Target::MemUnitFromKiB(long) @0x1d61bf20
GranuleTarget::ChunkGranules() @0x1d61a440 = (ChunkCellCount·4) / ChunkGranuleBytes
整芯片 HBM 速率Target::HbmFullChipBytesPerSecond() @0x1d6172a0（= Target+0x4f0）
每芯片核心数Target::CoresPerChip(TpuCoreType) @0x1d615b40
按层级 GB/s 矩阵Target::LocalDmaBandwidth(MemorySpace, MemorySpace) @0x1d6168e0（dispatcher）
复制模式消费者ShouldUseAsyncLocalCopy(HloInstruction, MemorySpace, MemorySpace, Target) @0x133eff40

两项成本 ​

目的 ​

每个已定价的操作数 DMA，即卷积 kernel 传输、reduce-window tile，或成本模型定价的任何 kCopy/kCopyStart，都会向 ResourceVector 放入两个周期值：一个延迟值进入启动 lane，一个带宽值进入传输 lane。两者分开放入，是因为调度器会独立地对 lane 取最大值；延迟 lane 类似 LLVM WriteLatency 的固定组成部分，带宽 lane 类似 getMemoryOpCost 的按大小成比例项。

算术 ​

c
// Per priced operand (input or output), inside RecordMemXferCyclesImpl:
granule        = Target.ChunkGranules()                            // @0x1d61a440
raw_bytes      = windowing_util::Size(wd.strides, MemUnit, granule)// align_up to granule
transfer_bytes = raw_bytes / (Compact2ndMinorRatio · ElementPackingFactor)

if rv[latency_lane] == 0.0:                                        // charged ONCE per lane
    rv.Acc(latency_lane, DefaultHbmInitLatency(MemUnit, wd, Target))   // @0x14552ca0
bytes_per_cycle = GetBytesPerCycle(hlo, Target, MemSpace)          // @0x1454dd00
rv.Acc(bw_lane, WindowCycles(MemUnit, wd, Target, bytes_per_cycle, 0, 0, 0))  // @0x14552660

rv[latency_lane] == 0.0 guard 是每 lane 一次规则：DMA 启动只在该 op 成本中第一次触碰某个 lane 时计费，因此进入同一 lane 的 N 次传输支付一次启动成本，而不是 N 次。带宽放入始终累加。输入/输出 lane 划分（输入为 R9/R10，输出为 R11/R12）属于 WindowDescription 字节成本，它为本页消耗的字节数定价。

注意 — 这两项被放入不同的 lane，并且在 bundle 级 MaxResourceCycles 归约之前永远不会加成一个数字。若一次传输的启动延迟（例如 v6e HBM 上 2280 周期）超过其带宽周期，它就是延迟受限：延迟 lane 主导 bundle max。大型连续传输则是带宽受限：bytes / bytes_per_cycle lane 主导。这种逐 lane 分离就是 roofline。

按层级启动延迟 — InitialDmaLatencyInNs ​

目的 ​

Target::InitialDmaLatencyInNs(MemorySpace dst, double burst_size) 返回开始一个 DMA 的固定纳秒启动成本，其中 DMA 的目标是给定层级。burst_size 参数会被接受，但每个 override 都会忽略它，也就是说延迟纯粹是按（dst 层级）计的常量，从不随传输大小变化。这是 DefaultHbmInitLatency 缩放到周期的值。

按层级常量 ​

从每个按代际 override 内部的 vmovsd xmm0, [rodata] 立即数解码而来（所有 .rodata double 均已按字节精确验证）：

dst MemorySpaceJellyfish (v2/v3)Pufferfish (v4)Viperfish (v5p)Ghostlite (v6e)base
VMEM (MS=3)24055500LogFatal
CMEM (MS=4)2405012001200LogFatal
SMEM (MS=5)24055512001200LogFatal
HBM / 任何其他24055512001200LogFatal

三种分派形态，每一种都已在反编译中按字节精确确认：

c
// JellyfishTarget @0x1d48f3a0 — unconditional constant, no MemorySpace branch
return 240.0;                              // vmovsd xmm0, [0xa2de6a8] = 240.0

// PufferfishTarget @0x1d493d00 — 2-entry table, index = (dst == kCmem)
idx = (dst == 4);                          // sete al
return ((double*)0xa2dcd40)[idx];          // [555.0, 50.0]  → CMEM=50, else=555

// ViperfishTarget @0x1d499ca0 / GhostliteTarget @0x1d496dc0 — VMEM-vs-other branch
return (dst == 3) ? 0.0 : 1200.0;          // vxorps→0; if dst!=VMEM, [0xa2df9f8] = 1200.0

// base Target @0x1d61c880 — LogMessageFatal("Unimplemented") at target.h:1235
LogFatal();                                // never reached for a configured target

延迟 rodata 常量（从 .rodata 解码）：

VAdouble作用
0xa2de6a8240.0Jellyfish，所有层级
0xa2dcd40555.0Pufferfish 默认值（表索引 0）
0xa2dcd4850.0Pufferfish CMEM（表索引 1）
0xa2df9f81200.0Viperfish / Ghostlite 非 VMEM

陷阱 — Pufferfish 的 CMEM 目标比其他每个层级都更便宜（50 ns vs 555 ns），这是唯一一个非 HBM 目标会降低启动成本的层级。0xa2dcd40 处的 2 项表由布尔值 dst == kCmem(4) 索引；如果重新实现硬编码单个 PF 延迟，就会把每个以 CMEM 为目标的 DMA 错估 11×。反过来，在 Viperfish/Ghostlite 上，唯一便宜的层级是 VMEM（0 ns），每个非 VMEM 目标，包括 CMEM 和 SMEM，都支付完整的 1200 ns。便宜层级的极性在 v4 和 v5p+ 之间翻转。

怪异点 — 基类 Target::InitialDmaLatencyInNs 是 target.h:1235 处的 LogMessageFatal("Unimplemented")，不是默认零。任何到达基类虚函数的 MemorySpace/代际组合都会中止编译。每个发布代际都覆盖它，因此这个 fatal 是防止未配置 target 的 guard，而不是活跃代码路径。

ns → 周期 — DefaultHbmInitLatency ​

目的 ​

DefaultHbmInitLatency 是成本模型中的适配器，它把按层级的纳秒常量转换为 TensorCore 周期，使启动延迟能与以周期计的带宽项和 CycleTable 项复合。它是放入延迟 lane 的值。

算法 ​

c
function DefaultHbmInitLatency(MemUnit mu, WindowDescription wd, Target t):  // @0x14552ca0
    sentinel = t.MemUnitFromKiB(0)                  // @0x1d61bf20 — the "no transfer" MemUnit
    if mu == sentinel: return 0.0                   // vxorpd — short-circuit, no DMA
    CHECK(mu.amount_ != INT64_MAX)                  // target.h:211 guard
    elem = (wd.shape != null) ? (WORD[wd.shape+0xb] >> 2) & 0x1f : 0   // operand element type
    init_ns = t.InitialDmaLatencyInNs(elem, window_bytes)             // vtable[+0x20]
    tc_ghz  = t.TensorCoreFrequencyInMegaHertz() / 1000.0             // /0xa2e0430 = /1000.0
    return init_ns * tc_ghz                                           // ns · GHz = cycles

该转换正是 ns · GHz：把 TC_MHz 除以 1000.0 得到 GHz，而 nanoseconds · gigahertz 是无量纲周期。/1000.0 除数已在 .rodata 0xa2e0430 = 1000.0 确认。对于 v6e 上的非 VMEM 目标（InitialDmaLatencyInNs = 1200 ns，Ghostlite TC = 1750 MHz）：1200 · 1.75 = 2100 周期会一次性放入延迟 lane。（1200 ns 是字节精确的；TC 频率不是 .rodata 常量，而是按 codename 的 chip_parts 字段，即 Target+0x90c，其值本身可恢复：v6e = 1750 MHz，v7x = 1900 MHz，两者都已对照 kDeviceTypeInfo tensorcore_clk_khz 表（+0x50，×1000 → Hz）以及按 codename 的 chip_parts Core[TC].frequency_mhz 行确认。）

陷阱 — 此处传给 InitialDmaLatencyInNs 的 MemorySpace 参数不是字面层级 enum，而是从 shape 指针重新提取的操作数元素类型字段（(WORD[shape+0xb] >> 2) & 0x1f）。对按层级延迟来说这目前无害，因为每个 override 要么忽略参数（JF），要么只分支少数值；但若未来某个 override 真的读取它，重新实现若传入层级 enum 而不是二进制传入的元素类型，就会分歧。

MemUnitFromKiB(0) 哨兵 ​

WindowCycles 和 DefaultHbmInitLatency 都先把传入的 MemUnit 与 Target::MemUnitFromKiB(0) 比较。MemUnit 是一个 {amount_, granule_tag} 对；MemUnitFromKiB(0) 是目标 granule 下规范的“零字节” MemUnit。当已定价访问是这个哨兵时，没有传输，两个函数都会立即返回 0.0，这是成本模型表示“该操作数位于不产生 DMA 的内存空间”的方式。

字节 → 周期 — MemUnitFromKiB 和 ChunkGranules ​

目的 ​

成本模型以 granule 单位而不是原始字节计数字节：传输在除以带宽之前总是向上取整到完整内存 granule。两个 accessor 定义 granule 和 KiB→granule 转换。

MemUnitFromKiB — @0x1d61bf20 ​

c
function MemUnitFromKiB(this, long amount_kib):                 // @0x1d61bf20
    granule_bytes = this.vtable[+0x5c0]()                       // ChunkGranuleBytes (per gen)
    CHECK(granule_bytes <= 1024)                                // target.cc:3244
    CHECK(1024 % granule_bytes == 0)                            // target.cc:3245 — power-of-two divisor
    shift = log2(1024 / granule_bytes)                          // _BitScanReverse
    return amount_kib << shift                                  // KiB → granule-units

MemUnitFromKiB(kib) 把 kibibyte 数转换为目标的原生 granule 单位：kib · (1024 / granule_bytes)，并以左移计算，因为 1024 / granule_bytes 被要求为 2 的幂（两个 CHECK 强制 granule_bytes | 1024 且 granule_bytes ≤ 1024）。当 amount_kib == 0 时结果为 0，即 DefaultHbmInitLatency/WindowCycles 测试的哨兵。

ChunkGranules — @0x1d61a440 ​

c
function ChunkGranules(this):                                  // @0x1d61a440
    chunk_cell_bytes = 4 * topology[+0x1a8]                     // ChunkCellCount · 4 bytes/cell
    granule_bytes    = this.vtable[+0x5c0]()                    // ChunkGranuleBytes (per gen)
    return chunk_cell_bytes / granule_bytes                     // granules per chunk

ChunkGranules 是 windowing_util::Size 把传输向上取整到的 granule：chunk-cell 字节大小（ChunkCellCount · 4）除以按代际的 ChunkGranuleBytes。Size 会对窗口的逐轴 stride count 求积，按该 granule 做 ceil-divide，再乘回去，从而把任何部分 granule 作为完整 granule 计费。MemUnitFromKiB 和 ChunkGranules 都读取 vtable[+0x5c0] 处同一个按代际的 ChunkGranuleBytes 虚函数，因此两个 accessor 的字节记账单位在代际内一致。

注意 — 取整 granule 是一个 chunk 概念，而不是 DMA descriptor granule。Size 的 align_up(Product(strides), ChunkGranules) 把字节体积向上取整到完整 chunk；DMA fragment 数（{1.0, 1.6, 1.3, 1.1, 1.05} 效率比）是在 WindowDescription 字节成本 中单独、正交的修正。把 chunk granule 与 descriptor count 混同的重新实现会重复修正。

字节/周期 — GetBytesPerCycle ​

目的 ​

bytes_per_cycle 是每核心、每周期的字节预算，WindowCycles 用它去除传输字节数。它是 roofline 的带宽分母，由整芯片字节速率和 TC 时钟推导得到。

算法 ​

c
function GetBytesPerCycle(HloInstruction* hlo, Target t, MemorySpace ms):  // @0x1454dd00
    CHECK(ms == kHbm(1) || ms == kCmem(4))               // fusion_util.cc:2354
    cycles_per_sec = t.TensorCoreFrequencyInMegaHertz() * 1e6     // *0xa2e0208 = 1e6
    full_chip_bps  = (ms == kHbm) ? t.HbmFullChipBytesPerSecond() // @0x1d6172a0 (Target+0x4f0)
                                  : t.CmemFullChipBytesPerSecond()// @0x1d6172c0
    bpc = full_chip_bps / cycles_per_sec / t.CoresPerChip(0)      // @0x1d615b40
    if ms == kHbm:
        env = GetTpuCompEnv(hlo.GetModule())[+0x1040]             // env HBM bw override
        if env > 0: bpc = env                                    // vblendvpd select
        // per-instruction backend_config FlagConfig override (opcode-14 custom-call path):
        if hlo has FlagConfig_Value of type kDouble: bpc = that.double_value  // [+0x10]
    return bpc

其几何关系是 bytes_per_cycle = FullChipBytesPerSecond / (TC_MHz · 1e6) / CoresPerChip：芯片的聚合内存速率先除以 Hz 时钟，转换为每周期字节，再分摊到芯片的 TensorCore 上。1e6 乘数已在 .rodata 0xa2e0208 = 1000000.0 确认；HbmFullChipBytesPerSecond 是 Target+0x4f0 处的普通字段读取（由 chip parts 填充）；CoresPerChip 从 topology 读取按 TpuCoreType 的计数。

来源地址含义
TensorCoreFrequencyInMegaHertz0x1d615b60TC 时钟（MHz）；纯 mov 0x90c(%rdi),%eax 读取 Target+0x90C 的 int32 字段，由 chip parts 填充（按 codename 的值已确认：v6e 1750，v7x 1900）
HbmFullChipBytesPerSecond0x1d6172a0Target+0x4f0 字段（整芯片 HBM B/s）
CmemFullChipBytesPerSecond0x1d6172c0整芯片 CMEM B/s（kCmem 路径）
CoresPerChip0x1d615b40来自 topology 的按 TpuCoreType 核心数
1e6 Hz 乘数0xa2e0208MHz → Hz

陷阱 — GetBytesPerCycle 只对 kHbm 和 kCmem 有效；任何其他 MemorySpace 都会在 fusion_util.cc:2354 触发 LogFatal("m_space == MemorySpace::kHbm || m_space == MemorySpace::kCmem")。VMEM/SMEM 传输不是通过这个函数定价的，它们的周期成本走 HBM bytes_per_cycle 分母以及 WindowCycles 内部按方向的 VMEM lane（如下），而不是 VMEM 专属的 bytes_per_cycle。两个覆盖出口，即 GetTpuCompEnv[+0x1040] 的芯片级 env 值和按 op 的 FlagConfig double，允许调用方为基准测试固定 bytes_per_cycle；忽略它们的重新实现会在 env 或 custom-call backend config 设置非零 HBM 带宽时分歧。

Roofline — WindowCycles ​

目的 ​

WindowCycles 是组装完整成本的公共入口：它加上启动延迟，用 bytes_per_cycle 除字节，并且在 v5p 及更新版本上分别为 HBM 和 VMEM 方向定价并取较慢者。它的返回值是带宽 lane 的周期值。按方向 max 与启动加法的组合就是 roofline。

算法 ​

c
function WindowCycles(MemUnit mu, wd, Target t, double bpc, long c, long hbm_b, long vmem_b):  // @0x14552660
    // source order is (mu, wd, t, bpc, c, hbm_b, vmem_b); bpc rides xmm0, the three longs r8/r9/stack
    if mu == t.MemUnitFromKiB(0): return 0.0           // sentinel — no transfer

    // startup-latency term (cycles)
    init = (c == -1) ? DefaultHbmInitLatency(mu, wd, t)             // @0x14552ca0
                     : (t.TensorCoreFrequencyInMegaHertz() / 1000.0) * c   // freq-scaled override

    count_desc = (t.vtable[+0x590]() == 1)             // per-gen "count descriptors" predicate
    base = WindowCyclesGenericTargetAgnostic(mu, wd, count_desc, bpc)  // @0x14552180 — the bytes/bpc divide
    cycles = base + init                               // VLOG-6 "original:" (fusion_util.cc:3435)

    if (hbm_b | vmem_b) == 0:                          // roofline OFF — the memory-xfer path (c=0)
        return cycles                                  // pure startup + bandwidth, single lane

    // roofline ON — only the conv classic-window path passes non-zero hbm_b/vmem_b
    CHECK(mu.amount_ != INT64_MAX)                     // target.h:211
    {hbm_lat, vmem_lat} = {t.InitialDmaLatencyInNs(kHbm=1,..), t.InitialDmaLatencyInNs(kVmem=3,..)}  // vtable[+0x20]
    if t[+0x398] < 5:                                  // < v5p (TpuVersion at Target+0x398)
        return (vmem_lat_scaled + hbm_lat_scaled) + max(dir0, dir1)   // single-lane shuffle/max
    else:                                              // v5p+ two-lane HBM/VMEM roofline
        // {hbm_bytes, vmem_bytes} / 1000.0 (xmmword_A2CE650), * TC_MHz, gated by >0.01
        hbm_cyc  = (hbm_b  / 1000.0) * TC_MHz
        vmem_cyc = (vmem_b / 1000.0) * TC_MHz
        return (vmem_lat_scaled + hbm_lat_scaled) + max(hbm_cyc, vmem_cyc)

核心除法 transferred_bytes / bytes_per_cycle 发生在 WindowCyclesGenericTargetAgnostic 内部（对传入 xmm0 的 bytes_per_cycle 执行 vdivsd），fragment-ratio 细节见 WindowDescription 字节成本。WindowCycles 总是在其上加启动 init；双 lane roofline 是一个由 (hbm_b | vmem_b) != 0 门控的可选块。当两个字节数都为零时（memory-xfer 成本路径会传 0, 0），WindowCycles 直接返回 base + init，并完全跳过 roofline 块。

双 lane 块（已在 @0x1455282e 确认）是显式的 roofline；在 v5p+（Target+0x398 >= 5）上，它独立计算每个方向：

  • 它加载 {hbm_bytes, vmem_bytes} 对，并把每个都除以 1000.0（xmmword_A2CE650 对，确认值为 [1000.0, 1000.0]），再乘以 TC 频率，把每个方向的字节数转换为各自的周期估计。
  • 对 xmmword_A2D8020（[0.01, 0.01]）执行 vcmpnlepd，以 1% 百分比下限门控每个方向，并通过 vandpd 选择该方向的延迟是否贡献。
  • vmaxsd 取按方向最大值，也就是 HBM 和 VMEM 中较慢者。这就是“双 lane roofline”：HBM 受限的传输和 VMEM 受限的传输会独立定价，而 op 支付更差者。

fusion_util.cc:3474 处的 VLOG-6 trace 命名了每个中间值，确认模型是双 lane max："hbm_percentage"、"vmem_percentage"、"hbm_bytes"、"vmem_bytes"、"mem_bw"、"vmem_byte_rate"、"hbmbw"、"vmembw"、"hbmlatency"、"vmemlatency" 和 "next_gen"（v5p+ 标记，fusion_util.cc:3481）。

rodata / vtable值 / 偏移作用
0xa2ce650[1000.0, 1000.0]按方向字节→速率除数（v5p+）
0xa2d8020[0.01, 0.01]按方向 1% 百分比门控（vcmpnlepd）
vtable[+0x590]== 1按代际的 count-descriptors 谓词
Target+0x398>= 5v5p+ 双 lane 门控（TpuVersion）

注意 — c（第 5 个）参数选择启动来源：c == -1 使用 DefaultHbmInitLatency；任何其他 c 使用 TC_GHz · c 作为显式的按调用启动计数（因此 c == 0 会在 WindowCycles 内部得到零启动）。RecordMemXferCyclesImpl 传入 c == 0，它会自己把按层级 DefaultHbmInitLatency 放入单独的延迟 lane（见§算术），并只把 WindowCycles 用于带宽项。传入 c == -1 的调用方是 permutation_util::IterateThroughWindowConfigs @0x145350c0 和 SelectAndScatterEmitter::EmitR3plus @0x10e71320，在那里启动被折叠进单个返回周期值。卷积 classic-window 路径（SpatialMajorConvolution::CalculateClassicWindowCost @0x131626c0）是唯一传入非零 hbm_b/vmem_b 的调用方，因此也是唯一到达双 lane roofline 块的调用方。

按层级带宽矩阵 — LocalDmaBandwidth ​

目的 ​

不同于 bytes_per_cycle 这个周期成本分母，按层级的 LocalDmaBandwidth<Src>To<Dst> 虚函数为每个有序（源层级，目标层级）对返回原始 GB/s。这些值供给复制模式决策，即一次复制应作为 async local DMA 发出，还是保持同步，而不是按字节周期成本。dispatcher Target::LocalDmaBandwidth(src, dst) @0x1d6168e0 把一个 (src, dst) 对映射到 15 个可达的按层级虚函数之一；每个都返回 std::optional<double>（IEEE-754 bits 在 rax，present flag 在 dl）。

矩阵（GB/s） ​

从每个按代际 override 中的 return <ieee754> 立即数解码；— = 无 override（基类返回 nullopt = 未建模）。Viperfish 列在 variant_name == "lite" 门控适用处显示 default / lite（所有立即数均已按字节精确验证）：

Pair (Src→Dst)DragonfishPufferfishViperfish (def/lite)Ghostlite
HBM → HBM—48072 / 30864
HBM → VMEM4234811198 / 8221285
HBM → SMEM—3455 / 5655
HBM → CMEM(via VmemToVmem slot)(via VmemToVmem)(via VmemToVmem)(via VmemToVmem)
VMEM → HBM42311111224 / 8281432
VMEM → VMEM—54472 / 82764
VMEM → CMEM—1121——
VMEM → SMEM—3455 / 5655
CMEM → HBM—1080——
CMEM → VMEM—2339——
CMEM → CMEM—1193——
CMEM → SMEM—34——
SMEM → HBM—3455 / 5655
SMEM → VMEM—3455 / 5655
SMEM → CMEM—34——
SMEM → SMEM—172828
SPMEM → HBM——587.4588

重新实现者必须遵守的结构事实：

  • 不存在 LocalDmaBandwidthHbmToCmem 虚函数。 dispatcher 把 (Hbm, Cmem) 路由到 VmemToVmem slot（vtable+0x1b8）：HBM↔CMEM 传输被成本模型视为经 VMEM staging，因此 HBM→CMEM 这一段按 VmemToVmem 速率定价。
  • 基类 Target 对每个 pair 都返回 0（nullopt），已在 Target::LocalDmaBandwidthHbmToVmem @0x1d492580 返回 0 处确认。缺失值表示“此层级对未建模”，ShouldUseAsyncLocalCopy 会把它视作“不要把它路由为 local DMA”，从而保持同步复制。
  • Dragonfish 只覆盖 HBM↔VMEM（二者都是 423，0x407A700000000000）；Jellyfish 没有定义自己的 LocalDmaBandwidth override，而是继承。
  • Viperfish 根据 "lite" 门控，通过 SSO-aware 的 variant_name.size() == 4 && variant_name == "lite" 比较（4 字节字面量 0x6574696C，'l','i','t','e'），已在 LocalDmaBandwidthHbmToVmem @0x1d49a380 确认：default → 1198，lite → 822。单值 VF pair（SmemToSmem = 28，SpmemToHbm = 587.4）不门控。

按字节精确确认示例：

c
DragonfishTarget::LocalDmaBandwidthHbmToVmem @0x1d48fa20 → 0x407A700000000000 = 423.0
GhostliteTarget::LocalDmaBandwidthHbmToVmem  @0x1d4973e0 → 0x4094140000000000 = 1285.0
ViperfishTarget::LocalDmaBandwidthHbmToVmem  @0x1d49a380 → "lite"? 822.0 : 1198.0
Target::LocalDmaBandwidthHbmToVmem (base)    @0x1d492580 → 0   (nullopt)

怪异点 — GB/s 矩阵和 bytes_per_cycle 预算是两个服务于两个不同决策的不同带宽数字。bytes_per_cycle（来自 GetBytesPerCycle，仅 HBM/CMEM）在成本模型决定发出传输后为该传输的周期成本定价。LocalDmaBandwidth<Src>To<Dst> 矩阵决定一次复制是否以 async 方式发出，并供给 collective-vs-local bisection-bandwidth 阈值。把 GB/s 矩阵用作周期成本分母（或反之）的重新实现会生成自洽但错误的成本模型。完整 dispatcher 映射和复制模式/bisection 消费者记录在 Local DMA Bandwidth。

调度器如何消费延迟/带宽 ​

延迟和带宽周期值落入 ResourceVector，并由 bundle 级 MaxResourceCycles 归约到 op 成本中。随后 CostModelLatencyEstimator（LatencyHidingScheduler 客户端）通过 cycles · trip_count / (TC_MHz · 1e6) 把周期转换为墙钟秒；这与 GetBytesPerCycle 使用的是同一套 TC 频率几何关系，因此带宽项从头到尾量纲一致。

复制模式决策位于此之前。ShouldUseAsyncLocalCopy @0x133eff40 读取 LocalDmaBandwidth(src, dst) GB/s 矩阵：如果 optional 缺失，复制保持同步；如果存在，对于多设备 collective，它会把本地速率（乘以 replicas² 并带一个 ×4 因子）与跨芯片 bisection 带宽比较，以选择 local-async 还是 collective-routed。该门控决定哪些 DMA 存在；本页的延迟/带宽周期模型为通过该门控的 DMA 定价。通用 XLA memory_space_assignment::CostAnalysis::GetAsyncCopyElapsed 使用单个平坦的配置带宽（DefaultMemBandwidthBytesPerSecond @0x1dceb320），不是这个按层级矩阵；按层级模型是 TPU 专属的，并在 MSA 的 kAlternate/kDefault 选择之前运行。

相关组件 ​

名称关系
WindowDescription 字节成本构造本模型除以 bytes_per_cycle 的字节数；拥有 R9–R12 lane 路由
Local DMA Bandwidth完整的 LocalDmaBandwidth(src,dst) dispatcher 映射 + 复制模式/bisection 消费者
成本模型概览三族成本模型以及 cycles → seconds TC 频率转换
CycleTable 家族带宽/延迟周期在 bundle max 中与之复合的吞吐整数
Resource Enum (23-slot)延迟和带宽值放入的 ResourceVector lane

交叉引用 ​

  • WindowDescription 字节成本 — windowing_util::Size、fragment ratio，以及消费此 bytes_per_cycle 和 DefaultHbmInitLatency 的 R9/R10/R11/R12 lane 划分
  • Local DMA Bandwidth — 按（src,dst）的 GB/s dispatcher @0x1d6168e0、HBM→CMEM-via-VmemToVmem 路由，以及 ShouldUseAsyncLocalCopy
  • 成本模型概览 — LatencyHidingScheduler 成本客户端，以及与 GetBytesPerCycle 共享的 TensorCoreFrequencyInMegaHertz 接线
  • CycleTable 家族 — 内存传输周期要与之取 max 的按代际吞吐后端
  • Resource Enum (23-slot) — 对延迟和带宽 lane 执行的 MaxResourceCycles 归约
  • Reduce-Window / Pooling 成本 — 通过此延迟/带宽模型为 operand-window DMA 定价的消费者
  • ConvolutionCostState — RecordConvKernelCycles，它把操作数 DMA 延迟/带宽与卷积计算一起放入