GTC：全局时钟与时钟同步树

第一章和本章前四节的计时全部用 LCC。LCC 数的是本芯片自己的时钟周期，这正是优化一段代码时要数的量；但它回答不了另外两类问题：

• 两个事件谁先谁后。 每个 TensorCore 有自己的 LCC，起点互不相关；不同芯片各有自己的时钟，快慢也不完全相同。两个 TensorCore 的 LCC 读数相减没有意义。
• 一段时间是多少秒。 LCC 的单位是“本芯片的一个周期”，换成秒要先知道这颗芯片此刻的时钟频率。

为此 TPU 另有一个计数器：全局时钟 GTC（global time counter）。它的设计目标是让整个切片的所有芯片读到同一条时间轴。本节先读出 GTC、看清它的读数由哪两部分组成，再说明各颗芯片的 GTC 是怎样对齐的（时钟同步树），以及由此带来的一个容易忽略的事实：作为时间源的芯片与其余芯片，读到的 GTC 行为不同。最后给出两个计数器的分工，并测出 GTC 在同一颗芯片、不同芯片之间的对齐精度。下一节的 XProf 就建立在 GTC 上。

读 GTC

GTC 与 LCC 的读法相同，也是 64 位，在同一个 bundle 中读低、高两半：

{ s0: srdreg.gtclo s21 ; s1: srdreg.gtchi s26 }

tpuasm_tools 中的 read_gtc(low) 生成这个 bundle，用法与第一章第 3 节的 read_lcc(low) 完全一样。两种读数都在标量一侧执行，不向 VIF 送入任何东西，所以相邻的 LCC 读数与 GTC 读数之间不需要 sfence（第 3 节）。

本节的片段要在一次运行中读 4 次，所以用 LccProbe.run_raw(body, reads)：它返回每次运行的全部 64 位读数，而不是 run 那样的两个差值。LccProbe(num_cores=2) 让一颗芯片的两个 TensorCore 执行同一段片段；进程打开多颗芯片时，LccProbe(device=k) 让载体在第 k 颗芯片上运行。

GTC 的读数由两部分组成

本小节实验源码、输出。实验在只打开一颗芯片的会话中进行。

每个周期加多少

最直接的问题：LCC 每个周期加 1，GTC 每个周期加多少？片段是连续 4 个读 GTC 的 bundle，没有等待时每个周期标量发射一个 bundle（第一章第 3 节），所以相邻两次读数恰好相隔 1 个周期：

body = read_gtc(20) + read_gtc(21) + read_gtc(22) + read_gtc(23)
counters = probe.run_raw(body, 4, repeats=REPEATS)[:, 0].astype(np.int64)

2000 次运行，相邻读数之差只出现三种组合：

(G1 − G0, G2 − G1, G3 − G2) 与次数：[((1, 15, 16), 651), ((15, 16, 1), 649), ((16, 1, 15), 700)]
G3 − G0：[32]
低 4 位的取值与次数：[(0, 5351), (1, 2649)]

每个周期的增量按 1、15、16 循环，三种组合只是从循环的不同位置开始读；任意连续 3 个周期合计 32。这个不均匀的增量不是噪声，它来自读数的结构。

高 60 位是时间，低 4 位是本地的序号

Google 的专利 US10365987B2《Synchronous hardware event collection》描述了这种计数器的一个实施例：64 位的 GTC 分成两部分，高 60 位是全局的时间计数，低 4 位是一个偏移（offset）。高位按全局时间推进；一个本地周期内高位没有变化时，低位加 1，使这个周期内发生的事件仍能与上一个周期区分；高位一推进，低位清零。

按这个结构把读数拆开，记高 60 位为 T = G >> 4，低 4 位为 O = G & 15：

增量 (1, 15, 16)：T = [0, 0, 1, 2]，O = [0, 1, 0, 0]
增量 (15, 16, 1)：T = [0, 1, 2, 2]，O = [1, 0, 0, 1]
增量 (16, 1, 15)：T = [0, 1, 1, 2]，O = [0, 0, 1, 0]

（T 以每次运行的第一次读数为 0。）三种组合是同一件事：每 3 个周期中，T 有 2 个周期各加 1，有 1 个周期不变；T 不变的那个周期 O 从 0 变成 1，下一个周期 T 加 1、O 回到 0。于是增量依次是：

低 4 位只出现 0 和 1，次数约为 2 : 1，也与此一致。

所以 1、15、16 的含义是：时间计数 T 每 3 个本地周期走 2 步。本地时钟约 1.05 GHz，T 的速率就是它的 2/3，即 700 MHz。libtpu 把 GTC 换算成时间时用的参数正是“频率 700 MHz、低 4 位另算”（研究报告 35）。把整个 64 位读数当作一个整数时，它平均每秒增加 700 MHz × 16 = 11.2 × 10⁹，这就是通常说的“GTC 是 11.2 GHz”；但芯片上并没有一个 11.2 GHz 的时钟，GTC 的一个计数也不是 1/11.2 ns 的均匀刻度。真正的时间刻度是 T 的一步，1/0.7 ≈ 1.43 ns，即 1.5 个周期。

长区间上的比例

要确认“3 个周期走 32”在长时间上也成立，片段用 LCC 包住 GTC，中间用 vdelay 让向量一侧停住 H 个周期，再用 sfence 让后面的读数等它结束：

setup = bundle(f's0: simm.s32 s24, {delay}')
body = read_lcc(20) + read_gtc(21) + bundle('misc: vdelay s24') + bundle('s0: sfence') + read_gtc(22) + read_lcc(23)

四次读数依次记作 L0、G0、G1、L1。实验只改 H，从 10³ 到 10⁹：

ΔL 总是 H + 15，完全确定：vdelay 的 H 个周期，加上 sfence 与几次读数的固定开销。GTC 的两次读数在 LCC 两次读数的内侧，两端各隔一个 bundle，所以 G0 到 G1 之间是 ΔL − 2 个周期。ΔG 在每一行都只取两三个值，彼此相差 1 或 15，这正是上面的三种相位；ΔG / (ΔL − 2) 随 H 增大收敛到 10.666667，即 32/3。

再把 H = 10⁸ 与 10⁹ 两次运行的主机时间相减（装载程序等固定开销在差值中抵消），得到以主机时钟为准的频率：

LCC 频率 1.0479 GHz，GTC 频率 11.178 GHz

主机计时的误差在千分之一左右，所以这里只能说“约 1.05 GHz、约 11.2 GHz”。研究报告 53 用约 33 秒、覆盖空闲、标量、DMA、向量和 MXU 负载的时间轴测得 GTC 与 11.2 × 10⁹ 相差约 4 ppm，且不随负载变化。前两章把周期换算成微秒时用的“约 1.05 GHz”就来自这里。

用时间轴标定 LCC 的频率

本小节实验源码、输出。

千分之一的误差来自主机：两次长运行各自的提交、唤醒时间并不完全相同。要把频率测得更准，可以利用 LCC 的一个性质：它在两次运行之间也持续计数。于是许多次短的运行可以串成一条长的时间轴，主机时钟的误差只出现在这条轴的两端，被整条轴的长度摊薄。

实验把第 1 节的 kernel（pl.delay(10000)）在约 20 秒内运行 200 次。每次运行，第一章第 3 节的 KernelClock 在 kernel 两端各读一次 LCC，主机在调用前后各读一次 CLOCK_MONOTONIC：

host_start = time.monotonic_ns()
timed(x).block_until_ready()
host_stop = time.monotonic_ns()
begin, finish = (int(value) for value in clock.read(2)[0])
samples.append((host_start, host_stop, begin, finish))
time.sleep(GAP)

第一次运行的起点读数发生在它的两次主机读数之间，最后一次运行的终点读数也夹在它的两次主机读数之间。所以整段 LCC 之差对应的真实时间，不短于“第一次的 host_stop 到最后一次的 host_start”，不长于“第一次的 host_start 到最后一次的 host_stop”：

cycles = last[3] - first[2]
shortest, longest = last[0] - first[1], last[1] - first[0]

第 1 轮：200 次运行，主机区间 20.257909 至 20.258379 秒，LCC 增加 21270904426 个周期
  频率的包络：1049.9806 至 1050.0049 MHz（宽度 23.2 ppm）；各次运行中点的拟合：1049.9927 MHz
第 2 轮：200 次运行，主机区间 20.031605 至 20.032131 秒，LCC 增加 21033354445 个周期
  频率的包络：1049.9809 至 1050.0084 MHz（宽度 26.3 ppm）；各次运行中点的拟合：1049.9933 MHz

两轮的包络都包含 1050 MHz，宽度二十几 ppm：LCC 的频率是 1050 MHz，与 TPU v4 公开的时钟频率一致（Jouppi 等，ISCA 2023）。这个包络不依赖提交和唤醒的延迟：延迟再大，也只是让两端的主机区间变宽。一个周期是 0.952 ns，1000 个周期是 0.95 µs。单芯片会话中 GTC 恰好是 LCC 的 32/3 倍，所以 GTC 的速率相应地是 11.2 × 10⁹，与研究报告 53 的结果一致。

反过来，不能用一次调用的主机耗时去除它的周期数。同一组数据中，一次调用的主机耗时约 240 µs，kernel 只有约 1.1 万个周期：

  单次调用的主机耗时：中位数 244.3 µs；若直接用一次调用的周期数除以它的主机耗时：46.7 MHz

主机耗时几乎全是提交与唤醒（第 1 节），这样算出的“频率”小了二十多倍。这种测法来自 top-k 论文的 3.3 节。

到这里为止，GTC 看起来只是 LCC 乘以 32/3。下面会看到，这只在一种情况下成立：这颗芯片自己就是时间的来源。

时钟同步树：各芯片的 GTC 怎样对齐

每颗芯片有自己的时钟。两颗芯片的时钟即使标称频率相同，实际也会相差百万分之几，而且各自从不同的时刻开始计数。要让所有芯片读到同一条时间轴，就得指定一颗芯片作为时间的来源，其余芯片不断地把自己的时间计数校正到它上面。

TPU v4 用一棵树做这件事。树的边就是芯片之间的 ICI 链路：

• 一颗芯片是时间源（global leader）。它的时间计数 T 直接由自己的时钟推进，不听从任何芯片。
• 其余每颗芯片是跟随者：它在自己的 6 个 ICI 端口中选定一个，从这个端口连着的那颗芯片（树上的父节点）接收时间，据此校正自己的 T；它又通过其余端口把时间传给自己的子节点。

专利中描述的做法是：父节点周期性地把自己的时间发给子节点；子节点比较收到的时间与本地的时间，并用一次往返（ping-and-echo）估计链路上的传输延迟，把它补偿掉，再据此加快或放慢自己的时间计数。TPU v4 具体的消息格式、校正周期和补偿公式没有公开，研究报告 53 也没有恢复出来；下面只使用能直接读回和测到的事实。

读回每个端口的角色

本小节实验源码、输出，在 TPU v4 (2x2x2) 的两个 host、八颗芯片上用 podrun 运行。

每个 ICI 端口有一组 GTC 寄存器，其中一个记录这个端口在同步树中的角色，一个是这个端口看到的 GTC 计数。它们可以从主机只读地读出来：以只读方式打开 /dev/accel{chip}，以只读方式映射寄存器所在的页，读一次 64 位。gtc_common.read_ports 做的就是这件事：

BAR2 = 0x10000000  # /sys/class/accel/accel*/bar_offsets 给出的 BAR2 在 /dev/accel* 中的映射起点
PORT_BASE = 0x37B80000  # 端口 0 的 GTC 寄存器组在 BAR2 中的偏移
PORT_STRIDE = 0x200000  # 相邻端口的寄存器组相隔这么多

page = mmap.mmap(descriptor, 4096, mmap.MAP_SHARED, mmap.PROT_READ, offset=BAR2 + PORT_BASE + port * PORT_STRIDE)
words = np.frombuffer(page, dtype=np.uint64)
# 偏移 0x00 是角色，0x28 是这个端口的 GTC 计数。
ports.append((int(words[0]), int(words[5])))

寄存器的位置和三种角色的取值由研究报告 53 从 libtpu 的配置代码中恢复：0 是 global leader（整棵树的时间源），1 是 local leader（这颗芯片从这个端口接收父节点的时间），2 是 follower（其余端口）。这个脚本没有写任何寄存器。

这里有一个容易出错的地方：/dev/accel0–/dev/accel3 的编号不是 JAX 的 device 编号，device.local_hardware_id 也不是它。JAX 的 device 没有给出这个编号的属性，所以 gtc_common.accel_index 实测对应关系：只在一个 device 上反复运行一个小程序，看哪颗芯片的中断数（/sys/class/accel/accel*/interrupt_counts）在增加。

before = interrupt_counts()
for _ in range(64):
    x = step(x)
x.block_until_ready()
increases = [after - earlier for after, earlier in zip(interrupt_counts(), before)]
chip, = [chip for chip, increase in enumerate(increases) if increase >= 64]

实验比较两个时刻：进程还没有 import jax 时，以及 jax.distributed.initialize() 之后。八颗芯片的结果（节选）：

device 0，坐标 (0, 0, 0)，进程 0，/dev/accel0
  会话之前：6 个端口的角色 ['follower']，计数 [15]
  会话之中：端口 0 follower、端口 1 global leader、端口 2 follower、端口 3 follower、端口 4 follower、端口 5 follower
  6 个端口的计数相隔 10 ms 两次读数都在增加：True；端口 1 增加 ...
device 1，坐标 (1, 0, 0)，进程 0，/dev/accel3
  会话之前：6 个端口的角色 ['follower']，计数 [15]
  会话之中：端口 0 follower、端口 1 local leader、端口 2 follower、端口 3 follower、端口 4 follower、端口 5 follower
device 2，坐标 (0, 1, 0)，进程 0，/dev/accel1
  会话之中：端口 0 local leader、端口 1 follower、端口 2 follower、端口 3 follower、端口 4 follower、端口 5 follower
...
device 7，坐标 (1, 1, 1)，进程 1，/dev/accel2
  会话之前：6 个端口的角色 ['follower']，计数 [15]
  会话之中：端口 0 local leader、端口 1 follower、端口 2 follower、端口 3 follower、端口 4 follower、端口 5 follower

从中读出四件事：

• device 与 /dev/accel* 的顺序不同。 两个 host 上都是 device 0、1、2、3（以及 4、5、6、7）依次对应 accel0、accel3、accel1、accel2。按 device 编号直接去读 /dev/accel*，会把三颗芯片的寄存器张冠李戴。
• 同步树是会话建立时配置的。 会话之前，48 个端口全是 follower，没有任何时间源；jax.distributed.initialize() 之后，整个切片恰好有一个 global leader 端口，在 device 0 上，其余七颗芯片各有恰好一个 local leader 端口。
• device 0 是时间源。 它是坐标 (0, 0, 0) 的芯片。
• 没有时间源时，计数停在 15。 会话之前每个端口读到的计数都是 15，即 T = 0、O = 15：时间计数没有推进，低 4 位从 0 数到 15 就停住，不再进位。这就是专利所说的“offset 数到最大值后饱和”，下一小节还会在运行中的芯片上看到它。

角色只说明每颗芯片从哪个端口接收时间，没有说明那个端口连着谁。研究报告 53 在同一个切片上进一步读出了每个端口对端的芯片与端口（按 /dev/accel* 的编号）：每个 host 上 accel1 与 accel3 的父节点是 accel0，accel2 的父节点是 accel1，第二个 host 的 accel0 的父节点是第一个 host 的 accel0。按上面的对应关系换成 device，就是完整的树：

device 0 (0,0,0)  时间源
├─ device 1 (1,0,0)
├─ device 2 (0,1,0)
│  └─ device 3 (1,1,0)
└─ device 4 (0,0,1)
   ├─ device 5 (1,0,1)
   └─ device 6 (0,1,1)
      └─ device 7 (1,1,1)

这棵树有一条简单的规则：每颗芯片的父节点，是把它的坐标中第一个不为 0 的分量减 1 得到的那颗芯片（按 x、y、z 的顺序）。所以一颗芯片在树上的深度就是它的坐标中不为 0 的分量个数，最深的 device 7 与时间源隔 3 条链路。这也是研究报告 53 从 libtpu 的建树代码中恢复出的规则。

时间源与跟随者读到的 GTC 不同

本小节实验源码、输出，同样在八颗芯片上用 podrun 运行。

现在在八颗芯片上各自重复本节开头的两个实验。相对上一个实验的片段没有任何改动，只是每个进程对自己的 4 颗芯片各建一个载体：

for index, device in enumerate(jax.local_devices()):
    probe = tpuasm_tools.LccProbe(device=index)

每颗芯片连续 4 次读 GTC，运行 4000 次；再停 10⁹ 个周期，比较 ΔG 与 32/3 × ΔL：

时间源完全规则。 device 0 的 4000 次运行全部是 1、15、16 的轮换，低 4 位只有 0 和 1；10⁹ 个周期上 3ΔG − 32ΔL 只取 −16 和 −13，也就是说 ΔG 与 32/3 × ΔL 之差始终不超过一个相位。它的 T 就是由自己的时钟每 3 个周期推 2 步，没有任何校正。

跟随者有约 1% 的读数偏离这个规律。 七颗跟随者各有 27–45 次运行读到别的组合，出现最多的几种是：

device 5：其他 35 次 [((1, 1, 1), 7), ((1, 31, 32), 5), ((32, 1, 31), 4), ((31, 32, 1), 4), ((1, 31, 16), 4), ((16, 1, 31), 2)]
device 7：其他 45 次 [((1, 1, 1), 17), ((32, 1, 31), 7), ((1, 31, 32), 3), ((31, 32, 1), 3), ((1, 15, 32), 2), ((31, 16, 1), 2)]

按“高 60 位是时间、低 4 位是序号”的结构，它们各自的含义是：

• (1, 1, 1)：连续几个周期 T 都不动，只有 O 在加 1。这是 T 在等待：本地的时间计数走到了父节点的前面，停下来等它。低 4 位因此出现 2、3、……直到 15，而时间源上最大只到 1。
• (1, 31, 32)、(31, 32, 1)、(32, 1, 31)：与规则的 1、15、16 相比，后两个增量各多 16，即 T 在一个周期里走了 2 步。这是 T 在追赶：本地的时间计数落在了父节点后面。
• 相邻读数完全相同：等待超过了 15 个周期时，O 数到 15 之后饱和，不再增加，而 T 还没有恢复推进，整个 64 位读数就停住不变。device 3、5、6、7 的低 4 位都出现过 15，并且各有 1 次运行读到了相邻两次完全相同的 GTC。这种情况很少（4000 次中 0–1 次，另一次运行中 device 7 有 5 次），但它说明 GTC 在跟随者上不是严格递增的。上一小节在会话之前读到的计数 15，是同一个现象的极端情形：T 始终不推进。

跟随者的 ΔG 与 ΔL 之比不是 32/3。 LCC 数的是本芯片时钟的周期，GTC 的 T 则被校正到时间源的时间上。两颗芯片的时钟频率有差别时，同样 10⁹ 个本地周期里，时间源走过的时间就不是恰好 32/3 × 10⁹ 个计数。device 1、3 的 ΔG 多出约 2.6 ppm，说明它们的时钟比时间源慢 2.6 ppm；device 4、6 少 3.5 ppm，device 5、7 少 1.5 ppm，它们的时钟比时间源快。上面看到的等待与追赶，就是在消除这个差别：时钟较快的芯片要不时让 T 停一停，较慢的要不时让 T 多走一步。实际观察到的校正是成批出现的（一次停住许多个周期，同一颗芯片上等待和追赶都会出现），这与“每隔一段时间收到父节点的时间、再一次性校正”的做法一致；校正的周期和算法本节没有恢复。

表中还有三个现象：

• 偏离量成对出现：device 1 与 3、4 与 6、5 与 7 几乎相同，device 2 与时间源 device 0 相同。每一对都是 x、z 坐标相同、y 坐标不同的两颗芯片，这与每两颗这样的芯片共用一个时钟源的情形一致；本节没有进一步验证时钟源的分配。
• 偏离量随时间漂移，成对的两颗一起漂移：同一天的三次运行中，device 1、3 依次是 +1.9、+1.8、+2.6 ppm，device 4、6 是 −4.1、−4.2、−3.5 ppm。时钟的频率本来就随温度变化，所以这些数只说明量级。
• device 2 的频率与时间源没有差别，却仍然有 34 次不规则的读数，与其他深度为 1 的芯片一样多。所以校正不只是为了消除频率差：即使两颗芯片的时钟一样快，跟随者的 T 也在被调整。它在补偿什么（例如对链路延迟估计的修正），本节没有查明。

为什么此前只看到 1、15、16

本节开头的实验，以及本教程中绝大多数实验，都在只打开一颗芯片的会话中运行。第一个实验的最后读出了这种会话中的角色：

端口 0 follower、端口 1 follower、端口 2 follower、端口 3 follower、端口 4 follower、端口 5 follower
相隔 10 ms 两次读数，6 个端口的计数都在增加：True

单芯片会话中没有同步树：没有任何端口被设为 leader，芯片不从任何端口接收时间，但 GTC 在推进，由芯片自己的时钟驱动。这颗芯片就是它自己的时间源，所以读到的总是规则的 1、15、16，ΔG 与 ΔL 之比恰好是 32/3。这与用的是 host 上的哪一颗芯片无关：分别只打开第 0、1、3 颗芯片，结果相同，而第 1、3 颗在八芯片会话中是跟随者。

因此，“GTC 每周期按 1、15、16 增加”“GTC 恰好是 LCC 的 32/3 倍”都不是 GTC 的一般性质，而是在时间源上测量的结果。程序在多芯片会话中运行时，除 device 0 之外的每颗芯片都是跟随者：GTC 的相邻读数之差可以是 0，可以是 31、32，一段区间的 ΔG 也不能用 32/3 精确地换成本地周期数。

两个计数器的分工

由此得到使用规则：

• 在一个 TensorCore 上数一段代码要多少周期，用 LCC。 优化一个 kernel、验证发射模型、比较同一颗芯片上的两种写法，都属于这一类。读数就是周期，没有换算，也不受同步校正的影响。本教程中的周期数都是这样测的。
• 比较不同 TensorCore、不同芯片上的事件，用 GTC。 例如多芯片 kernel 中各芯片何时开始、何时结束。LCC 在这里不能用：各 TensorCore 的 LCC 起点不同，各芯片的周期长短也不完全相同。
• 把一段时间换成秒，用 GTC。 ΔG / 11.2 是纳秒数。只在足够长的区间上这样做：短区间上 ΔG 受相位和校正的影响，时间源上有 ±1 步（约 1.4 ns），跟随者上一次校正可以有十几步。只打开一颗芯片时，LCC 的周期数除以 1.05 GHz 给出同样的结果。
• 需要区分先后时不要只靠 GTC。 跟随者上两个不同的事件可以读到相同的 GTC；同一个 TensorCore 内的先后用 LCC 或程序自身的顺序判断。

同一颗芯片的两个 TensorCore

GTC 被设计成整个切片共用的时间轴，但两个 TensorCore 的读数实际上对得多准，要测了才知道。困难在于：一方读一次、另一方读一次，无从知道两次读取相隔多久。两个 TensorCore 之间唯一能传递“先后”的手段是发信号，而信号在路上要花时间。下面先用最保守的办法给出上下界，再利用“两个 TensorCore 共用一个时钟”这一条件把偏移精确地算出来。

只用因果顺序：得到上下界

本小节实验源码、输出，在只打开一颗芯片的会话中运行。

每个 TensorCore 执行同一段片段：

BODY = (
    read_gtc(20)                                                 # G0
    + bundle('misc: vsyncadd.remote.s32 [sflag:s29], 1')         # 给对方发信号
    + bundle('misc: vwait.ge [sflag:45], 1')                     # 等对方的信号
    + bundle('misc: vsyncadd.s32 [sflag:45], -1')
    + bundle('s0: sfence')
    + read_gtc(21)                                               # G1
    + read_lcc(22)
)

s29 是对方 TensorCore 的 sflag 45 的地址，在 setup 中按编译器生成入口汇合的同样方式拼出（第二章第 1 节）。sflag 45 正是入口汇合用的同步标志，片段中加 1、减 1 各一次，不影响之后的出口汇合。

TensorCore 0 先读 G0、再发信号；TensorCore 1 收到信号、经过 sfence 之后才读 G1。所以 TensorCore 0 的 G0 在真实时间上早于 TensorCore 1 的 G1，反方向同理。记 o 为同一时刻 TensorCore 1 的 GTC 减去 TensorCore 0 的 GTC，两条因果关系给出：

G0[1] − G1[0] ≤ o ≤ G1[1] − G0[0]

一次运行给出的区间宽约 960–3460 个 GTC 计数，就是信号在两个 TensorCore 之间往返一次的时间。32 次运行的区间取交集：

32 个区间的交集：[-481, 481] 个 GTC 计数，即 [-42.9, 42.9] ns

这个结论不依赖任何假设，但很弱：它只说明偏移在 ±43 ns 之内。区间的宽度是信号往返的时间，再多做几次也不会变窄。原因是这个办法对信号的时延一无所知：往返一共花了约 90 个周期，但去程和回程各占多少不知道，“去程 1 个周期、回程 89 个周期、偏移 +44 个周期”与“去程回程各 45 个周期、偏移为 0”给出完全相同的读数。要得到更准的结果，必须加入关于时延的信息。

利用共用的时钟：精确到一个计数

本小节实验源码、输出。

同一颗芯片的两个 TensorCore 有一个可以利用的条件：两个 LCC 以完全相同的速率计数。也就是说，在任何时刻，TensorCore 1 的 LCC 减去 TensorCore 0 的 LCC 都是同一个常数，记作 D。这一条件不是硬件文档给出的，而是要由实验确认的前提，下面第一步的数据同时检验它。D 不知道，但它是一个整数，而且不随时间变化。只要求出 D，就能把一个 TensorCore 上的任何 LCC 读数换成另一个 TensorCore 上同一时刻的 LCC 读数，两边的 GTC 读数也就能换到同一时刻比较。

片段相对上一个实验只多了两次 LCC 读数，分别紧挨着发信号之前和收到信号之后：

BODY = (
    read_gtc(20)
    + read_lcc(21)                                               # 发出时刻
    + bundle('misc: vsyncadd.remote.s32 [sflag:s29], 1')
    + bundle('misc: vwait.ge [sflag:45], 1')
    + bundle('misc: vsyncadd.s32 [sflag:45], -1')
    + bundle('s0: sfence')
    + read_lcc(22)                                               # 收到时刻
    + read_gtc(23)
)

第一步：求 D。 TensorCore 0 发出信号时自己的 LCC 是 sent[0]，TensorCore 1 收到时自己的 LCC 是 received[1]。信号在路上花了 d₀₁ 个周期，而两个读数来自相差 D 的两个计数器，所以：

received[1] − sent[0] = d₀₁ + D        TensorCore 0 → 1
received[0] − sent[1] = d₁₀ − D        TensorCore 1 → 0

两式相加消去 D，得到往返时延 d₀₁ + d₁₀；相减得到 2D + (d₀₁ − d₁₀)。如果两个方向的时延相等，D 就求出来了。

时延每次运行并不相同：信号到达时，对方可能还没有执行到 vwait，那么“收到时刻”就被推迟了。但时延有一个下限：对方已经在 vwait 上等着，信号一到就继续。这个下限只取决于硬件和片段本身。两个 TensorCore 是相同的硬件，执行的是同一段片段，所以两个方向的最小时延应当相等。400 次运行中每个方向最小的几个值：

TensorCore 0 → 1：收到时刻 − 发出时刻，最小的几个值：-51253（206 次）、-51252（21 次）、-51251（35 次）、-51250（6 次）
TensorCore 1 → 0：收到时刻 − 发出时刻，最小的几个值：51345（382 次）、51346（1 次）、51349（1 次）、51352（2 次）
两个最小值之和 92，之差 -102598：单程时延 d = 46 个周期，LCC 之差 D = -51299

两个方向都有一个清晰的下限，而且大多数运行就落在下限上，说明最小时延确实是确定的。往返 92 个周期，每个方向 46 个周期（其中包括 vsyncadd.remote 的发射、vwait 被唤醒、sfence 的 11 个周期等片段自身的开销），D = −51299：TensorCore 1 的 LCC 比 TensorCore 0 的小 51299。

这组数据同时检验了“D 是常数”这个前提。如果两个 LCC 的速率有哪怕极小的差别，D 就会随时间漂移，下限也跟着移动，同一个最小值不可能在整个实验期间反复出现。实际上：

第一次与最后一次运行相隔 141818016 个周期；0 → 1 取到最小值的运行从第 0 次到第 397 次，各四分之一中的次数 [45, 55, 52, 54]
1 → 0 取到最小值的运行从第 0 次到第 399 次，各四分之一中的次数 [98, 98, 92, 94]

400 次运行跨越 1.4 × 10⁸ 个周期，两个方向的最小值从头到尾都是同一个整数，在四段时间里出现得同样频繁。D 在这段时间里的漂移不到 1 个周期，即两个 LCC 的速率之差小于 10⁻⁸（百分之一 ppm）。对比之下，不同芯片的时钟相差百万分之几（上文的表），同样长的时间里会漂移几百个周期。这与两个 TensorCore 的 LCC 由同一个时钟驱动一致，而不像是两个恰好接近的时钟；本节没有电路层面的依据，只确认了算法需要的“D 是常数”。

第二步：把 GTC 换到同一时刻。 上面看到，时间源上 GTC 每 3 个周期恰好走 32，所以在一个 TensorCore 上，3G − 32L（G、L 是同一时刻的 GTC 与 LCC）只随 GTC 的三种相位取三个值。每个 TensorCore 上各有两对相邻的读数，算出来正是如此：

TensorCore 0 的 3G − 32L：[-23447325501, -23447325488, -23447325472]
TensorCore 1 的 3G − 32L：[-23445683933, -23445683920, -23445683904]

在同一时刻，TensorCore 0 的 LCC 是 L，TensorCore 1 的 LCC 就是 L + D；两边处在同一种相位时，两个 GTC 之差 o 满足 3o = (3G − 32L)[1] − (3G − 32L)[0] + 32D。三种相位分别代入：

同一时刻两个 TensorCore 的 GTC 之差的 3 倍，按三种相位分别计算：[0, 0, 0]

偏移恰好是 0。 同一颗芯片的两个 TensorCore 在同一时刻读到的是同一个 GTC 值。对比只用因果顺序得到的 [−512, 512]，结论从“相差不超过 46 ns”变成了“相等”。

这个结论用了一条假设：两个方向的最小时延相等。它不能由这个实验本身证明，但数据对它有很强的约束。o 必须是整数，所以 32D 加上两边 3G − 32L 的差必须是 3 的倍数，D 只能取 −51299 + 3k。k ≠ 0 意味着两个方向的最小时延相差 6 个周期以上，而且偏移恰好是 32 个计数的整数倍；在相同的硬件上执行相同的片段，没有理由出现这种差别。

由这个实验还得到两个副产品：

• 两个 TensorCore 的 LCC 可以互相换算。 第一章第 3 节说不同 TensorCore 的 LCC 不能直接相减，原因只是不知道 D。D 测出来之后（本机上是 −51299，几次会话中都相同，但硬件不保证它；上一个实验中紧随 G1 的两个 LCC 读数相差 −51329 到 −51278，就是 D 加上两边到达时刻的差），同一颗芯片两个 TensorCore 上的事件就可以用 LCC 精确到周期地比较先后。
• 谁先到、差多少。 在上一个实验中，两个 TensorCore 读 G0 的时刻相差 −224 到 2784 个计数，即 −20 到 249 ns：它们都在入口汇合之后、输入 DMA 之后到达同一位置，到达的先后和差距每次运行都不同。第二章第 1 节所说的“谁先到、差多少”，就可以在汇合前后各插一次这样的读数测出。

不同芯片之间

本小节实验源码、输出，在 TPU v4 (2x2x2) 的八颗芯片上用 podrun 运行。

跨芯片时，GTC 的对齐精度取决于同步树。上一小节的办法在这里只能用一半：不同芯片的时钟不同，LCC 之差不再是常数，不能借 LCC 求出精确的换算。但“取两个方向的最小时延、假设它们相等”仍然可以直接用在 GTC 读数上。

写法

kernel 在八颗芯片上运行。一颗芯片作为发起方，给应答方发一个信号；应答方一直在等，收到后立即回信：

@pl.when(me == initiator)
def _() -> None:
    # 先停 20 µs，让应答方已经在等待；再发信号，等回信。
    pl.delay(20000)
    pl.semaphore_signal(ping, 1, device_id={'device': responder, 'tc': 0})
    pl.semaphore_wait(ping, 1)
    # 只用来标出这个分支的末尾，供插入读数时定位。
    pl.delay(8)

@pl.when(me == responder)
def _() -> None:
    pl.semaphore_wait(ping, 1)
    pl.semaphore_signal(ping, 1, device_id={'device': initiator, 'tc': 0})

ping 是一个普通信号量（pltpu.SemaphoreType.REGULAR，第二章第 4 节）。这与同一颗芯片上的实验不同：那里两边同时发、同时等，这里是一问一答。一问一答时，每个方向的接收方都已经在 vwait 上等着（应答方从一开始就在等，发起方发完立即等），所以每次运行测到的都是没有额外等待的时延。

然后用 tpuasm 在四个位置插入读 GTC 的 clock_read（第一章第 3 节的 KernelClock 用的同一段，只是把 lcc 换成 gtc），读数留在各自芯片的 SMEM 中：

# 发起方：发信号之前读 A0，等到回信之后读 A3。应答方：等到信号之后读 B1，回信之前读 B2。
timed = tpuasm_tools.load(tpuasm_tools.insert_bundles(serialized, {send: gtc(0), delays[1]: gtc(1), reply - 1: gtc(0), reply: gtc(1)}), compiled, jax.devices())

两个方向上，“读数 → 发信号”和“被唤醒 → 读数”都是同样的指令序列，所以读数自身的开销在两个方向上相同。程序跨越两个 host 时，改写后的 executable 要在每个进程中装载，并把全部 device 交给 load。每次运行之后，每个进程用自己的 KernelClock(device=k) 取回本 host 四颗芯片的读数，最后用 jax.make_array_from_process_local_data 和 process_allgather 把两个 host 的读数合到一起。

记 o 为应答方的 GTC 减去发起方的 GTC，d 为单程时延：

B1 − A0 = d(去程) + o
A3 − B2 = d(回程) − o

只用因果顺序（d > 0）时，o 在 −(A3 − B2) 与 B1 − A0 之间；假设两个方向的最小时延相等时，o = (min(B1 − A0) − min(A3 − B2)) / 2。为了排除发起与应答两种角色本身的差别，device 0 与每颗芯片 t 做两遍，分别由 device 0 和 t 发起，各 300 次。

结果

下表的 o 都是 t 的 GTC 减去 device 0 的 GTC，单位是 GTC 计数：

能证明的界。 只用因果顺序时，相邻的芯片是约 ±6400 个计数（±0.57 µs），隔两条链路的是约 ±11300（±1.0 µs），最远的 device 7 是约 ±16000（±1.4 µs）。这些界就是信号的单程时延：一条链路约 0.57 µs，每多一条链路多约 0.43 µs。它们说明“偏移不超过这么多”，不说明偏移实际有这么大。研究报告 53 用另一种写法得到了同样量级的界。

假设对称时的估计。 七颗芯片两遍的平均都是负的，在 −8 到 −356 个计数之间，即落后 device 0 约 1–32 ns，比上面的界小一到两个数量级。两种发起方式得到的符号和大小一致，说明这不是发起与应答两种角色的差别造成的。估计值大致随树上的深度增加：深度为 1 的三颗芯片是 −8、−28、−116，深度为 2 的是 −140、−220、−232，深度为 3 的 device 7 是 −356。这与“每条边补偿传输延迟之后留下一点偏差、沿路径累加”的预期方向一致。

这个估计有多可靠。 比同一颗芯片上的结果差得多，原因有三个：

• 最小值没有收敛。 同一颗芯片上，400 次运行中有一半以上恰好落在最小值上；这里 300 次运行中只有个别几次接近最小值，中位数比最小值大 250–700 个计数。两种发起方式的估计相差最多 64 个计数（device 5）；不同时间的运行之间差别更大：另一次没有留档的运行中，device 1 的平均是 −80，device 4 是 −24，与上表相差约 100 个计数。所以单颗芯片的估计值有约 ±100 个计数（±9 ns）的不确定，可靠的只是整体的量级和随深度增大的趋势。
• 两个方向不保证对称。 相邻芯片的去程和回程走同一条链路的两个方向；隔两条链路时，按先 x 后 y 的路由（第二章第 7 节），去程和回程经过的中间芯片不同。估计值中“t 落后”与“回程比去程慢”无法区分。
• 偏移不是常数。 跟随者的 GTC 在不断被校正，一次校正可以有十几步。

所以这个实验的结论是：各芯片的 GTC 实际对齐在几十纳秒的量级，深度越大偏得越多；能够不加假设地证明的，仍然只是微秒量级的界。

同时要知道这些数字背后与同步树有关的三件事：

• 误差沿树累积。 每条边在补偿传输延迟之后都可能留下一点偏差，一颗芯片相对时间源的偏差是从根到它的路径上各边偏差之和。深度为 3 的 device 7 经过三次这样的补偿。研究报告 53 在 0.5 ms 的窗口上比较了 3ΔG − 32ΔL 的抖动：时间源约 18，深度为 1 的芯片约 160–180，深度为 2 的约 260–310，device 7 约 400，随深度增大。
• 偏移随时间变化。 上面的估计把偏移当作固定值来求，实际的偏移在一个小范围内变化。
• 启动时没有验收精度。 研究报告 53 发现，TPU v4 的驱动在建立同步树之后只检查计数已经开始推进，没有检查各芯片的误差是否小于某个阈值。

跨芯片比较时间时，可以把 GTC 读数直接相减。若需要严格的保证，按微秒量级的精度解读：它足以比较两颗芯片上一次 all-reduce 的开始和结束（第二章第 7 节，几到几十微秒），不足以比较单条 remote DMA 的到达时刻（约 2 µs）。若接受上面的估计，实际的偏差在几十纳秒以内。第二章第 5 节提到的不同芯片之间的时间比较，就要按这个精度解读。