Completion Loop 和 AsyncTrackingEvent
本页中的所有地址都适用于
libtpu-0.0.40-cp314wheel 中的libtpu.so（buildlibtpu_lts_20260413_b_RC00，build-id md589edbbe81c5b328a958fe628a9f2207d）。该镜像未 strip；demangled C++ 符号名按原样引用。.textVMA 等于文件偏移。其他版本会有所不同。
摘要
当 PJRT 调用方启动一个程序时，libtpu 不会把一个可轮询的 stream handle 交给 host。它返回一个 future，即包装 tsl::AsyncValue 的 tsl::Future<void>，并且 runtime 承诺在设备 retire 该 launch 时让该 async value available。整个完成模型建立在 tsl::AsyncValue 上：一个引用计数、单赋值的 cell，保存“not-yet”、一个值或一个错误，并在转换为 available 的瞬间触发已入队的 waiter 列表。这是 TFRT（TensorFlow Runtime）async-value 模型，也是现代 PJRT 路径：xla::TpuClient 经由 tpu::System，并作为 PjRtFuture / PJRT_Event async 系列暴露给 C API。它不是 Stream Semantics 记录的旧版 stream_executor::Stream + RecordEvent/WaitFor 机制。不存在 host poll loop 在设备 flag 上旋转；完成是 push 模式，在 AsyncValue fulfilled 时作为 waiter callback 交付。
一次 launch 串起三个关联对象。（1）enqueue 路径通过 xla::CommonPjRtClient::CreateLinkedUserPromise（0xf918980）生成一个链接的 promise 对：一半是 tsl::PromiseBase<absl::Status>，其 Future 返回给用户；另一半是设备侧 xla::PjRtDeviceEventPromise（一个 TpuTrackedDeviceEventPromise），包装 tsl::AsyncValueRef<tpu::TpuEvent>。两者被链接起来，因此解析设备事件会转发到用户 future。（2）该设备事件注册为 launch 的 define event：xla::TpuHostTransferManager::SetExecuteEvent（0xf813760）把它盖到 host transfer manager 上，tpu::System::Execute（0x1d0b33e0）在其 define_events span 中接收它。（3）当程序在设备上完成时，tpu::System::Execute 的 completion lambda 运行 tpu::TpuEventIssuer::FulfillArgs，解析 define event；随后 TpuTrackedDeviceEventPromise::SetReady（0xf837920）ForwardTo 用户的 async value，触发所有 waiter，包括 output-buffer releaser 和用户的 OnReady done-callback。
本页负责completion-event 生命周期：per-launch event 如何创建（CreateLinkedUserPromise / CreateLinkedEventPromise / CreateAsyncTrackingEvent）、如何作为 define event 注册并携带进入 tpu::System::Execute、设备侧 TpuEventIssuer sequence-point 引擎如何在 launch retire 时 fulfil 它，以及 fulfilment 如何扇出到用户 done-callback（PJRT_Event_OnReady）和 output-buffer release。产生 launch 的 enqueue 位于 LoadProgramAndEnqueueToStream；旧版 SE 路径上的 cross-stream WaitFor/RecordEvent wait primitives 属于 Stream Semantics（链接即可，不要重复）；execute 入口是 ExecuteAsyncOnStream。
对重新实现而言，契约是：
tsl::AsyncValue完成原语 — 一个引用计数、单赋值 cell，其 state byte 门控“available”，其& 8allocated flag 和_InterlockedDecrement驱动销毁，且 waiter list 是唯一通知通道。没有 polling。- 链接的 promise 对 —
CreateLinkedUserPromise返回{ PromiseBase<absl::Status> user_promise, AsyncValue* device_event }；用户获得tsl::Future，runtime 保留设备侧 promise，并把两者链接起来。 - define-event 注册 + fulfilment 路径 —
SetExecuteEvent注册 per-launchTpuEventdefine event；tpu::System::Execute在define_events中携带它；TpuEventIssuer::{NextSequencePoint, RunWhenDepsReady, FulfillArgs}按设备完成顺序 fulfil 它；TpuTrackedDeviceEventPromise::SetReady通过tsl::IndirectAsyncValue::ForwardTo将其转发到用户值。 - done-callback 扇出 — fulfilment 触发
AsyncValuewaiter list：通过PJRT_Event_OnReady注册的 callback 会运行，输出缓冲区的 definition event 变为 available，因此其后备 HBM 会被释放。
| 完成原语 | tsl::AsyncValue（引用计数单赋值；& 8 = allocated，_InterlockedDecrement destroy） |
| 面向用户的句柄 | tsl::Future<void> → PJRT_Event（槽位 10–14） |
| 链接的 promise 对 | xla::CommonPjRtClient::CreateLinkedUserPromise @ 0xf918980（776 B） |
| 设备事件 promise | xla::TpuTrackedDeviceEventPromise（Set/SetReady/SetError）基于 tsl::AsyncValueRef<tpu::TpuEvent> |
| 按设备事件工厂 | xla::TpuDevice::CreateAsyncTrackingEvent @ 0xf7ff1c0（891 B）· C-ABI PJRT_Device_CreateAsyncTrackingEvent |
| define-event 注册 | xla::TpuHostTransferManager::SetExecuteEvent @ 0xf813760（570 B） |
| launch + define events | tpu::System::Execute @ 0x1d0b33e0（10370 B），define_events = 最后一个 span 参数 |
| 设备 sequencer | tpu::TpuEventIssuer::{NextSequencePoint 0x1d0d38e0, Sequence::Next 0x1d0d3940, RunWhenDepsReady 0x1d0d4640, FulfillArgs} |
| Resolve → forward | TpuTrackedDeviceEventPromise::SetReady @ 0xf837920 → tsl::IndirectAsyncValue::ForwardTo |
| Done-callback 槽位 | pjrt::PJRT_Event_OnReady @ 0xf86fc60 · PJRT_Event_Await @ 0xf86fa80 · PJRT_Event_IsReady @ 0xf86f9e0 |
| 证据等级 | 重新实现级 / 对 IDA 反编译字节确认 |
1. 完成原语 — tsl::AsyncValue
目的
本页的一切都归结为一个对象：tsl::AsyncValue。它是一个 heap cell，引用计数，生命开始时 unavailable，并且只会被赋值一次，赋为一个值（tpu::TpuEvent、absl::Status、void）或一个错误。赋值是触发其 waiter list 的唯一边。runtime 从不在 host 上轮询 device-resident completion flag；相反，host 在 async value 上注册 callback，device-completion lambda 会 fulfil 该值并运行 callback。这种反转区分了现代 PJRT 路径和 Stream Semantics 中旧版 SE BlockHostUntilDone/RecordEvent poll-and-notify 模型。
布局和 Refcount/State 纪律
completion 路径上触及的每个 AsyncValue 都遵循相同字节纪律，在 CreateLinkedUserPromise、SetReady 和 tpu::System::Execute 的 waiter node 中可见相同模式：
// The AsyncValue refcount + state idiom (observed at every release site)
function AsyncValue_DropRef(av): // e.g. SetReady 0xf837920 lines 36-54
if av == nullptr: return
if (av->byte[4] & 8) == 0: // bit 3 of +4 = "heap-allocated"
return // statically-owned: never destroy
if av->refcount == 1 || // +0x00 = atomic refcount (uint32)
AtomicDecrement(&av->refcount) == 0:
AsyncValue::Destroy(av) // last ref: free
```text
三个对重新实现关键的事实：（1）偏移 `+0` 处的 32 位 word 是 atomic refcount；release 使用 `_InterlockedDecrement`，并且当 count 已经为 1 时有跳过 atomic 的 fast-path。（2）偏移 `+4` 处的 byte 携带 flags；**bit 3（`& 8`）表示“heap-allocated”**。静态拥有的 async value（进程 singleton ready event）可通过该 bit 清零识别，并且*永不*销毁，所以同一段 drop-ref 代码可安全用于两者。（3）在把值交给 logger 或 waiter list 之前 acquire ref，是在同一个 `& 8` 测试保护下执行 `_InterlockedIncrement`（`SetReady` line 26）。释放非 heap async value 或用 polling 代替注册 waiter 的重新实现都会破坏模型。
> **注意 —** 具体 waiter-list 机制（`tsl::AsyncValue::EnqueueWaiter<Lambda>` 为每个调用生成带 `RunWaiterAndDeleteWaiterNode` 的 `Node` 类型）会为每种 closure type 生成一次；二进制在 runtime 各处携带了每个 closure 的不同实例化，包括 `tpu::System::Execute` 和 `SetExecuteEvent` 下的 `EnqueueWaiter<…::$_0>` waiters（每个都有自己的 typeinfo、vtable、`Node` ctor/dtor 和 `RunWaiterAndDeleteWaiterNode`）。形状始终相同：一个持有 moved closure 的 node，在 `WaitersAndState` CAS 下串入该值的 waiter list，并在 fulfilment 时运行后删除。
---
## 2. 创建每次执行事件 — 链接的 Promise 对
### 目的
enqueue 路径需要同一完成事件的*两个*视图：返回给用户的 `tsl::Future<void>`（以便 JAX/PyTorch-XLA 可以 `await` 或附加 done-callback），以及 runtime 保留的设备侧 promise，使其能在 launch retire 时解析事件。`xla::CommonPjRtClient::CreateLinkedUserPromise` 一次性生成两者并把它们*链接*起来，因此一次解析会从设备传播到用户。
### 算法
去掉样板 `__function::__policy` move/destroy 噪声后，`CreateLinkedUserPromise`（`0xf918980`）的反编译在字节层面很清晰：
```c
// xla::CommonPjRtClient::CreateLinkedUserPromise(memspace, file, line, label) 0xf918980
// returns a 0x50-byte struct: { user_promise, device_event, profiling_callbacks... }
function CreateLinkedUserPromise(out, client, memspace, file, line, label):
PromiseMaker<void>::Make(&user_promise, &linked_av, profiling_keys) // line 41
// user_promise : tsl::PromiseBase<...>; linked_av : the AsyncValue the user's Future observes
// (vtable +536) CreateLinkedEventPromise(memspace, label) — mint the DEVICE-side promise
client->vtable[+536](&device_event_av, memspace, file, line, label, &linked_av) // line 77
// (vtable +528) CreateUserFuture / TrackFuture — wrap linked_av as the user Future
client->vtable[+528](client, memspace, file_b, line_b, &device_event_av) // line 87
out[0] = user_promise // PromiseBase<absl::Status> the runtime fulfils
out[1] = device_event_av // AsyncValue the device-side promise resolves
// out[0x10..0x40] = profiling key callbacks (TraceMe span open/close)
return out这两个 vtable 调用就是链接点。槽位 +536 是虚函数 CreateLinkedEventPromise；对 TPU client 而言，它是 xla::TpuClient::CreateLinkedEventPromise（0xf7faae0，610 B），会基于新的 tsl::AsyncValueRef<tpu::TpuEvent> 构造设备侧 TpuTrackedDeviceEventPromise，并把它接线到刚创建的 linked_av。槽位 +528 生成用户可见 future（TrackFuture/profiled-future 包装器，§5）。返回结构体的前两个 qword 是 runtime 将 fulfil 的 promise，以及用户将 await 的 async value，也就是同一个完成事件的两个句柄。
特性 —
CreateLinkedEventPromise是虚函数，并由每个 client 覆盖：CommonPjRtClient（基类）、TpuClient（TPU 设备事件）、PjRtCpuClient（CPU），并且对拒绝它的 client 有 guard 字符串"CreateLinkedEventPromise is not supported"。重新实现必须通过该 client 的 override 路由；链接语义（哪个AsyncValueforward 到哪个）位于那里，而不是 common framework 中。
TpuTrackedDeviceEventPromise — 设备半边
设备侧 promise 是 xla::TpuTrackedDeviceEventPromise，即为 TPU 特化的 xla::PjRtDeviceEventPromise。它的表面由三个 setter 加一个 accessor 构成：
| 方法 | 地址 | 效果 |
|---|---|---|
Set(PjRtDeviceEventRef) | 0xf837a40 | 绑定具体设备事件（例如 transfer 中的 TpuEvent） |
SetReady() | 0xf837920 | 解析为 ready（tpu::ReadyTpuEvent）事件并 forward 到用户值 |
SetError(absl::Status) | 0xf83c580 | 将该值解析为错误状态（poisons future） |
async_value() | 0xf83c560 | 底层 tsl::AsyncValue*（用于 waiter 注册） |
SetReady 是标准解析路径，也是 forward 步骤最清楚的字节证据。其反编译（0xf837920）：
// xla::TpuTrackedDeviceEventPromise::SetReady() 0xf837920
function SetReady(this):
client = this->client_via_device->...() // this[2] -> device -> client (vtable +16) line 18
loc = this->device->shared_mem_location() // vtable +24 line 19
logger = client->pending_event_logger(loc) // optional event tracer
ready = tpu::ReadyTpuEvent() // an already-available TpuEvent line 21
if logger:
logger->Log(ready, " TpuTrackedDeviceEventPromise::SetReady", 39) // diagnostic string line 27
indirect = this->indirect_av // this[3] = tsl::IndirectAsyncValue* line 32
IndirectAsyncValue::ForwardTo(indirect) // <-- splice ready event into user value line 35
// drop refs on ready + indirect per the §1 discipline
```text
机制是 `tsl::IndirectAsyncValue::ForwardTo`：promise 持有一个*间接* async value（之前返回给用户的 placeholder），`ForwardTo` 会把现在 ready 的具体 `TpuEvent` 拼接进去。间接值变为 available；其 waiter list 被触发。`SetError` 对错误状态执行对称操作；`Set` 则 forward 一个调用方提供的设备事件，而不是 ready 事件。这是设备完成变为 host 可见的单一点。
---
## 3. CreateAsyncTrackingEvent — 具名按设备事件
### 目的
`CreateLinkedUserPromise` 是 execute 路径内部使用的进程内生成器。PJRT C-API 还暴露了一个*独立*事件工厂，使 framework 可以创建 tracking event、跨边界传递并稍后解析它，也就是本页标题的显式对应物。这是 `xla::TpuDevice::CreateAsyncTrackingEvent`（`0xf7ff1c0`，891 B），通过 `pjrt::PJRT_Device_CreateAsyncTrackingEvent` 从 C-ABI 到达。
### 函数映射
| 函数 | 地址 | 角色 |
|---|---:|---|
| `pjrt::PJRT_Device_CreateAsyncTrackingEvent` | (slot) | C-ABI 入口；解包 `PJRT_Device_CreateAsyncTrackingEvent_Args` |
| `xla::TpuDevice::CreateAsyncTrackingEvent(string_view label)` | `0xf7ff1c0` | 生成带 label 的设备 tracking event（返回 promise/future 对） |
| `xla::PjRtCpuDevice::CreateAsyncTrackingEvent` | — | CPU sibling（相同表面，CPU events） |
| `xla::MegaScalePjRtDevice::CreateAsyncTrackingEvent` | `0xe6eb780` | multi-slice decorator；转发到被包装设备 |
`string_view` 参数是人类可读 label（与 `CreateLinkedUserPromise` 的 `file`/`line`/`label` 三元组中传递的 label 相同），供 `SetReady` line 27 中看到的 `pending_event_logger` 诊断路径使用。tracking-event 调用生成的事件是同一种基于 `tsl::AsyncValue` 的对象；调用方通过返回的 promise（`SetReady`/`SetError`）解析它，方式与 execute 路径内部完全相同。
> **注意 —** 按设备工厂和 per-launch `CreateLinkedUserPromise` 生成的是*同一种*对象，即链接的 `{promise, AsyncValue}` 对。差异只在*谁*解析它：tracking event 由创建它的 framework 解析（例如用来门控 dependent launch），而 execute-path event 由 `tpu::System::Execute` 的 completion lambda 解析（§4）。两者都作为 `tsl::Future` / `PJRT_Event` 暴露给用户。
---
## 4. 注册并 Fulfill Define Event
### 目的
生成 event 只是生命周期的一半；另一半是把它*绑定*到 launch 上，使设备在 retirement 时解析它。per-launch event 注册为 **define event**（它*定义* launch 完成和输出缓冲区 readiness），携带进入 `tpu::System::Execute`，并在程序 drain 时由设备侧 `TpuEventIssuer` fulfil。
### SetExecuteEvent — 盖上 Define Event
在 `xla::TpuExecutableLoadState::ExecuteLaunchRaw`（`0xf8109a0`，per-launch driver）内部，`CreateLinkedUserPromise` 生成 completion pair 后，设备侧 event 会注册到 host transfer manager：
```c
// xla::TpuHostTransferManager::SetExecuteEvent(AsyncValueRef<tpu::TpuEvent>&) 0xf813760
function SetExecuteEvent(this, execute_event):
// store the launch's TpuEvent as THE execute (define) event for this transfer manager,
// and register a waiter so host-transfer completion is gated on it.
this->execute_event = execute_event // retain (AcquireRef)
AsyncValue::EnqueueWaiter(execute_event.av, on_execute_ready_lambda, ...) // $_0 waiterfunctions JSON 确认了该 waiter：tsl::AsyncValue::EnqueueWaiter<...SetExecuteEvent(...)::$_0>。SetExecuteEvent 把一个 continuation 串到 execute event 的 async value 上，使 host-side transfer bookkeeping（以及 host transfer manager 自身 teardown）在 launch 完成时运行。这就是 output device→host copy 得知计算完成的方式。
tpu::System::Execute — 携带 define_events
tpu::System::Execute（0x1d0b33e0，10370 B）是 device-runtime 入口。其签名是 wait/define event 模型的决定性证据（从 demangled symbol 字节确认）：
// tpu::System::Execute(...) 0x1d0b33e0
tpu::System::Execute(
AsyncValueRef<tpu::ProgramHandle> program,
api::ExecuteOptions options,
Span<AsyncValueRef<tpu::TpuBufferBase>> inputs,
Span<AsyncValueRef<tpu::TpuBufferBase>> outputs,
Span<AsyncValueRef<tpu::TpuEvent>> wait_events, // launch BLOCKS until these are ready
Span<AsyncValueRef<tpu::TpuEvent>> define_events) // launch FULFILS these on completion
```text
`wait_events` 是输入依赖集合（程序在这些 event 解析前不会启动，包括 input H2D copy、之前的 launch）。`define_events` 是输出集合，也就是*这个* launch 在 retire 时使其 available 的 event。来自 `CreateLinkedUserPromise` 的 per-launch event 位于 `define_events` 中。完成边是一段注册在 `Execute` 内部的 lambda（`$_0`），它在设备完成时运行并调用 `tpu::TpuEventIssuer::FulfillArgs`：
```c
// the completion lambda registered inside tpu::System::Execute Execute::$_1
function on_device_complete(fulfill_args): // FulfillArgs, run by the TpuEventIssuer
// success path:
issuer.Fulfill(fulfill_args) -> { // resolves each define_event
for ev in define_events:
ev.promise.SetReady() // §2: ForwardTo user value
}
// error path uses the sibling $_3 lambda taking FulfillOnErrorArgs ->
// ev.promise.SetError(status) // poisons the futuresTpuEventIssuer — 设备 Sequencer
旧版 SE 路径用 per-stream FIFO（DeepseaRequestQueue，见 Stream Semantics）排序工作。现代路径使用 tpu::TpuEventIssuer：一个 sequence-point + dependency-DAG 引擎，完全用 tsl::AsyncValue 依赖表达 launch 的 FulfillArgs lambda 何时可以运行。
| 函数 | 地址 | 角色 |
|---|---|---|
tpu::TpuEventIssuer::NextSequencePoint(int) | 0x1d0d38e0 | 分配下一个 sequence point（排序 token） |
tpu::TpuEventIssuer::Sequence::Next() | 0x1d0d3940 | 推进 sequence（链接连续 launch） |
tpu::TpuEventIssuer::RunWhenDepsReady(...) | 0x1d0d4640 | 注册 fulfil/fulfil-on-error callbacks，在所有 deps 的 AsyncValue 解析后运行 |
tpu::TpuEventIssuer::AggregateDeps<Span<...>>(...) | — | 将 input buffer/event refs 折叠为一个 dependency vector |
tpu::CreateCountTrackingTpuEventIssuer(int, ConcurrentWorkQueue*, SystemEventTracker*) | 0x1d0d3820 | 计数 in-flight events 的 issuer 变体（tracking） |
tpu::CreateNonTrackingTpuEventIssuer(int, ConcurrentWorkQueue*) | 0x1d0d37a0 | 轻量 issuer 变体（无计数） |
RunWhenDepsReady 是核心：其签名接收 define event（AsyncValueRef<TpuEvent>）、两个 RCReference<AsyncValue> 依赖 span、一个 FulfillArgs 成功 callback、一个 FulfillOnErrorArgs 错误 callback，以及一个 SequencePoint async value。它注册 callback，使其在所有依赖 available 后触发，这正是 SE 通过 WaitFor 达成的 DAG join。当 join 完成时，成功 callback 运行 FulfillArgs（解析 define event → SetReady → 用户值），或错误 callback 运行 FulfillOnErrorArgs（→ SetError）。
特性 — 存在两种 issuer 变体，选择会改变完成 bookkeeping，而非排序。
CountTrackingEventIssuer（用SystemEventTracker*构造）会在 event issue 和 fulfil 时增减 in-flight count，用于 runtime 必须报告 pending work 的场景（例如 throttling、telemetry）。NonTrackingEventIssuer跳过该计数。两者实现相同的IssueEventImpl(IssueArgs)虚函数；硬编码其中一种的重新实现要么丢失 in-flight accounting，要么在不需要时为其付费。
5. Done-Callback 扇出和 Buffer Release
目的
define event 的 fulfilment 是原因；触发 async value 的 waiter list 是结果。该列表上有两个不同消费者：用户 done-callback（通过 PJRT_Event_OnReady 注册）和输出缓冲区 release 路径（它们的 definition event 变为 available，允许在所有引用释放后释放其 HBM）。
用户 Future 和 PJRT_Event
链接对的用户半边由 xla::TpuClient::TrackFuture（0xf7fad60，258 B）包装成 tsl::Future<void>；当启用 profiling 时则由 CreateProfiledFuture（0xf7fae80，445 B）包装，它用 tsl::FutureHelpers::ProfilingKeys 键控的 TraceMe span open/close callback 包裹 future。该 future 被装箱进 PJRT_Event（C-API 槽位 10–14）。done-callback 表面如下：
| C-API 槽位 | 函数 | 地址 | 效果 |
|---|---|---|---|
PJRT_Event_IsReady | pjrt::PJRT_Event_IsReady | 0xf86f9e0 | 非阻塞：async value 是否 available？ |
PJRT_Event_Await | pjrt::PJRT_Event_Await | 0xf86fa80 | 阻塞调用方直到 available（唯一阻塞路径） |
PJRT_Event_OnReady | pjrt::PJRT_Event_OnReady | 0xf86fc60 | 注册在 fulfilment 时运行的 callback，即 push done-callback |
PJRT_Event_Error | pjrt::PJRT_Event_Error | 0xf86fba0 | 如果该值解析为错误，则读取 error status |
PJRT_Event_OnReady 是标准完成交付路径：它在底层 async value 上入队一个 waiter。当 SetReady 的 ForwardTo 使该值 available（§2）时，waiter 会在 runtime 线程上运行用户 callback。PJRT_Event_Await 是阻塞 fallback，是 host 线程在完成上停驻的唯一位置；大多数调用方会避免它，改用 OnReady。
Buffer Release
从 ExecuteLaunchRaw 返回的输出缓冲区在调用返回时还不可用；每个缓冲区都携带自己的 definition event（同一个 define event，或从它派生的 event）。缓冲区后备 HBM 在其 definition event 解析前不能复用。由于 define event 就是 launch-completion event，fulfilment（§4）会同时：（1）使输出缓冲区可读（其 definition event available），以及（2）释放 runtime 在 launch 持续期间保持存活的引用。当 TpuBuffer 的最后一个 RCReference 释放时（用户 release 加上 runtime 现在已触发的内部 hold），该 tpu::TpuBuffer 会返回给分配器。错误走平行路径：SetError 会把用户 future 和输出 definition event 解析为 error status，因此依赖方会看到失败，而不是读取未初始化的设备内存。
易错点 — 用户 future 解析为“ready”意味着 launch retired，并不意味着 device→host 输出 copy 已完成。输出缓冲区驻留在设备上；在 host 上读取它们仍需要
CopyRawDeviceToHostAndReturnEvent，其自身的TpuEvent受该完成门控。把 execute future 当作“输出已在 host memory 中”的重新实现会读取设备句柄。execute future 门控设备侧有效性；copy event 门控host 侧可用性。
错误传播和 Ready Shortcut
重新实现者必须镜像两个边界行为。（1）xla::TpuClient::CreateErrorEvent(absl::Status, PjRtMemorySpace*)（0xf808420，493 B）会生成一个已出错 event，用于 launch 在到达设备前即被拒绝的情况（坏参数、poisoned executable）。它会立即解析为错误，因此依赖方快速失败。（2）tpu::ReadyTpuEvent()（用于 SetReady 内部）返回一个已 available 的 TpuEvent，它是进程共享 singleton，可通过清零的 & 8 allocated bit（§1）识别，因此对它 drop ref 是 no-op。ready shortcut 意味着没有真实设备等待的 launch（或平凡完成的 launch）无需往返即可解析其 define event。
6. 重新实现注意事项
- Push，而不是 poll。 完成通过触发 async value 的 waiter list 交付，绝不通过 host 在设备 flag 上旋转。注册
OnReadywaiter；把Await留给少见的阻塞调用方。poll loop 不符合该 runtime 的形状。 - 链接对是一个完成事件、两个句柄。 不要把用户 future 和设备 event 建模为必须保持同步的独立对象。
CreateLinkedUserPromise一次性链接它们（ForwardTo）；解析设备侧会自动解析用户侧。双向同步逻辑是这里不存在的 bug 面。 - 尊重
& 8allocated bit。 同一段 drop-ref/destroy 代码会同时作用于 heap async value 和进程共享 ready/error singleton。销毁非 heap 值（& 8清零）会破坏 singleton。每次Destroy都必须由该 bit 门控。 - define_events 和 wait_events 不能互换。
wait_events阻塞 launch；define_events由 launch fulfil。交换它们会导致死锁（launch 等待它本应产生的 event）或永远不发出完成信号。 - 有意选择 issuer 变体。
CountTracking维护 in-flight accounting（telemetry/throttle）；NonTracking不维护。两者的排序 DAG（RunWhenDepsReady）相同，只有 bookkeeping 不同。 - Execute-ready ≠ outputs-in-host-memory。 execute future 门控设备侧有效性；单独的 D2H copy event 门控 host 可读性（§5 易错点）。把两者都暴露给调用方，否则他们会像对待 host pointer 一样读取设备句柄。
- 这仅是现代路径。 这里不涉及
stream_executor::Stream::RecordEvent/WaitFor/BlockHostUntilDone。旧版 SE completion 模型（absl::Notification支撑的 host events、device-queue waiters）是 Stream Semantics 记录的并行栈；PJRT launch 不会进入它。
相关组件
| 名称 | 关系 |
|---|---|
xla::CommonPjRtClient::CreateLinkedUserPromise | 生成 enqueue 路径返回的链接 {user_promise, device_event} 对 |
xla::TpuTrackedDeviceEventPromise | 该对的设备半边；SetReady/Set/SetError 解析它并 forward 到用户值 |
xla::TpuDevice::CreateAsyncTrackingEvent | 独立的按设备 tracking-event 工厂；进程内生成器的 PJRT C-API 对应物 |
xla::TpuHostTransferManager::SetExecuteEvent | 注册 per-launch define event，使 host transfer 受完成门控 |
tpu::System::Execute | device-runtime launch；携带 wait_events / define_events 并注册 completion lambda |
tpu::TpuEventIssuer | 设备 sequence-point + dependency-DAG 引擎；RunWhenDepsReady 在完成时触发 FulfillArgs |
tsl::AsyncValue / tsl::Future<void> / PJRT_Event | 完成原语、其面向用户的句柄，以及 done-callback 通过其注册的 C-ABI event |
交叉引用
- Overview — 本 completion 生命周期收尾的 PJRT 到设备 execute 路径
- ExecuteAsyncOnStream — 开始一次 launch 的执行入口，本页追踪其完成
- LoadProgramAndEnqueueToStream — 产生 launch 及其 per-execution define event 的 enqueue
- Stream Semantics — 旧版 SE stream 模型；负责
WaitFor/RecordEvent/BlockHostUntilDone（本页是现代 async-value 替代路径，不是重复） - Host Callbacks —
DoHostCallbackWithStatus，其absl::Status通过 completion event 流回 - PJRT Events & Async — 本生命周期驱动的 C-ABI
PJRT_Event/PjRtFuture表面（槽位 10–14） - PJRT Client & Device — 生成链接 promise 对的
xla::TpuClient/CommonPjRtClient - PJRT Buffer & Memory — fulfilment 扇出时其 definition event 变为 available 的输出缓冲区